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

[ir] Spit builder_impl tests. 

This Cl splits apart the builder_impl test file in order to make it
easier to determine what has been tested.

Bug: tint:1718
Change-Id: Ib2f4d11a398b26fd774eb8c60555475053f20e0b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/131581
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
This commit is contained in:
dan sinclair 2023-05-04 18:01:26 +00:00 committed by Dawn LUCI CQ
parent 34f41c7bad
commit 63716c55cc
10 changed files with 1113 additions and 894 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

54
builder_impl_builtin_test.cc Normal file
View File

@ -0,0 +1,54 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/ir/test_helper.h"
#include "gmock/gmock.h"
#include "src/tint/ast/case_selector.h"
#include "src/tint/ast/int_literal_expression.h"
#include "src/tint/constant/scalar.h"
namespace tint::ir {
namespace {
using namespace tint::number_suffixes; // NOLINT
using IR_BuilderImplTest = TestHelper;
TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) {
auto i = GlobalVar("i", builtin::AddressSpace::kPrivate, Expr(1_f));
auto* expr = Call("asin", i);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, f32, read_write> = var private read_write
store %1:ref<private, f32, read_write>, 1.0f
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
%2:f32 = asin %1:ref<private, f32, read_write>
ret
func_end
)");
}
} // namespace
} // namespace tint::ir

179
builder_impl_function_test.cc Normal file
View File

@ -0,0 +1,179 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/ir/test_helper.h"
#include "gmock/gmock.h"
#include "src/tint/ast/case_selector.h"
#include "src/tint/ast/int_literal_expression.h"
#include "src/tint/constant/scalar.h"
namespace tint::ir {
namespace {
using namespace tint::number_suffixes; // NOLINT
using IR_BuilderImplTest = TestHelper;
TEST_F(IR_BuilderImplTest, EmitFunction_Vertex) {
Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(0_f, 0_f, 0_f, 0_f))},
utils::Vector{Stage(ast::PipelineStage::kVertex)},
utils::Vector{Builtin(builtin::BuiltinValue::kPosition)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@vertex ra: @position]
%fn1 = block
ret vec4<f32> 0.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_Fragment) {
Func("test", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{Stage(ast::PipelineStage::kFragment)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():void [@fragment]
%fn1 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_Compute) {
Func("test", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{Stage(ast::PipelineStage::kCompute), WorkgroupSize(8_i, 4_i, 2_i)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():void [@compute @workgroup_size(8, 4, 2)]
%fn1 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_Return) {
Func("test", utils::Empty, ty.vec3<f32>(), utils::Vector{Return(vec3<f32>(0_f, 0_f, 0_f))},
utils::Empty);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec3<f32>
%fn1 = block
ret vec3<f32> 0.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPosition) {
Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(1_f, 2_f, 3_f, 4_f))},
utils::Vector{Stage(ast::PipelineStage::kVertex)},
utils::Vector{Builtin(builtin::BuiltinValue::kPosition)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@vertex ra: @position]
%fn1 = block
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPositionInvariant) {
Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(1_f, 2_f, 3_f, 4_f))},
utils::Vector{Stage(ast::PipelineStage::kVertex)},
utils::Vector{Builtin(builtin::BuiltinValue::kPosition), Invariant()});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@vertex ra: @position @invariant]
%fn1 = block
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation) {
Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(1_f, 2_f, 3_f, 4_f))},
utils::Vector{Stage(ast::PipelineStage::kFragment)}, utils::Vector{Location(1_i)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@fragment ra: @location(1)]
%fn1 = block
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnFragDepth) {
Func("test", utils::Empty, ty.f32(), utils::Vector{Return(1_f)},
utils::Vector{Stage(ast::PipelineStage::kFragment)},
utils::Vector{Builtin(builtin::BuiltinValue::kFragDepth)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():f32 [@fragment ra: @frag_depth]
%fn1 = block
ret 1.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnSampleMask) {
Func("test", utils::Empty, ty.u32(), utils::Vector{Return(1_u)},
utils::Vector{Stage(ast::PipelineStage::kFragment)},
utils::Vector{Builtin(builtin::BuiltinValue::kSampleMask)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():u32 [@fragment ra: @sample_mask]
%fn1 = block
ret 1u
func_end
)");
}
} // namespace
} // namespace tint::ir

5
src/tint/BUILD.gn
View File

@ -2142,7 +2142,12 @@ if (tint_build_unittests) {
sources = [
"ir/binary_test.cc",
"ir/bitcast_test.cc",
"ir/builder_impl_binary_test.cc",
"ir/builder_impl_call_test.cc",
"ir/builder_impl_literal_test.cc",
"ir/builder_impl_materialize_test.cc",
"ir/builder_impl_test.cc",
"ir/builder_impl_var_test.cc",
"ir/constant_test.cc",
"ir/discard_test.cc",
"ir/store_test.cc",

5
src/tint/CMakeLists.txt
View File

@ -1440,7 +1440,12 @@ if(TINT_BUILD_TESTS)
list(APPEND TINT_TEST_SRCS
ir/binary_test.cc
ir/bitcast_test.cc
ir/builder_impl_binary_test.cc
ir/builder_impl_call_test.cc
ir/builder_impl_literal_test.cc
ir/builder_impl_materialize_test.cc
ir/builder_impl_test.cc
ir/builder_impl_var_test.cc
ir/constant_test.cc
ir/discard_test.cc
ir/store_test.cc

456
src/tint/ir/builder_impl_binary_test.cc Normal file
View File

@ -0,0 +1,456 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/ir/test_helper.h"
#include "gmock/gmock.h"
#include "src/tint/ast/case_selector.h"
#include "src/tint/ast/int_literal_expression.h"
#include "src/tint/constant/scalar.h"
namespace tint::ir {
namespace {
using namespace tint::number_suffixes; // NOLINT
using IR_BuilderImplTest = TestHelper;
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Add) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Add(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = add %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Subtract) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Sub(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = sub %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Multiply) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Mul(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = mul %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Div) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Div(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = div %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Modulo) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Mod(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = mod %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_And) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = And(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = and %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Or) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Or(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = or %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Xor) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Xor(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = xor %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
Func("my_func", utils::Empty, ty.bool_(), utils::Vector{Return(true)});
auto* expr = LogicalAnd(Call("my_func"), false);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():bool
%fn1 = block
ret true
func_end
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1:bool = call my_func
%2:bool = var function read_write
store %2:bool, %1:bool
branch %fn4
%fn4 = if %1:bool [t: %fn5, f: %fn6, m: %fn7]
# true branch
%fn5 = block
store %2:bool, false
branch %fn7
# if merge
%fn7 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
Func("my_func", utils::Empty, ty.bool_(), utils::Vector{Return(true)});
auto* expr = LogicalOr(Call("my_func"), true);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():bool
%fn1 = block
ret true
func_end
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1:bool = call my_func
%2:bool = var function read_write
store %2:bool, %1:bool
branch %fn4
%fn4 = if %1:bool [t: %fn5, f: %fn6, m: %fn7]
# true branch
# false branch
%fn6 = block
store %2:bool, true
branch %fn7
# if merge
%fn7 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Equal) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Equal(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = eq %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_NotEqual) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = NotEqual(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = neq %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThan) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = LessThan(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = lt %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThan) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = GreaterThan(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = gt %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThanEqual) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = LessThanEqual(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = lte %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThanEqual) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = GreaterThanEqual(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = gte %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftLeft) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Shl(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = shiftl %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftRight) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Shr(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = shiftr %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
Func("my_func", utils::Empty, ty.f32(), utils::Vector{Return(0_f)});
auto* expr = LogicalAnd(LessThan(Call("my_func"), 2_f),
GreaterThan(2.5_f, Div(Call("my_func"), Mul(2.3_f, Call("my_func")))));
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():f32
%fn1 = block
ret 0.0f
func_end
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1:f32 = call my_func
%2:bool = lt %1:f32, 2.0f
%3:bool = var function read_write
store %3:bool, %2:bool
branch %fn4
%fn4 = if %2:bool [t: %fn5, f: %fn6, m: %fn7]
# true branch
%fn5 = block
%4:f32 = call my_func
%5:f32 = call my_func
%6:f32 = mul 2.29999995231628417969f, %5:f32
%7:f32 = div %4:f32, %6:f32
%8:bool = gt 2.5f, %7:f32
store %3:bool, %8:bool
branch %fn7
# if merge
%fn7 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) {
Func("my_func", utils::Vector{Param("p", ty.bool_())}, ty.bool_(), utils::Vector{Return(true)});
auto* expr = Call("my_func", LogicalAnd(LessThan(2.4_f, 2_f),
GreaterThan(2.5_f, Div(10_f, Mul(2.3_f, 9.4_f)))));
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():bool
%fn1 = block
ret true
func_end
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1:bool = call my_func, false
ret
func_end
)");
}
} // namespace
} // namespace tint::ir

153
src/tint/ir/builder_impl_call_test.cc Normal file
View File

@ -0,0 +1,153 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/ir/test_helper.h"
#include "gmock/gmock.h"
#include "src/tint/ast/case_selector.h"
#include "src/tint/ast/int_literal_expression.h"
#include "src/tint/constant/scalar.h"
namespace tint::ir {
namespace {
using namespace tint::number_suffixes; // NOLINT
using IR_BuilderImplTest = TestHelper;
TEST_F(IR_BuilderImplTest, EmitExpression_Bitcast) {
Func("my_func", utils::Empty, ty.f32(), utils::Vector{Return(0_f)});
auto* expr = Bitcast<f32>(Call("my_func"));
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:f32 = call my_func
%2:f32 = bitcast %1:f32
)");
}
TEST_F(IR_BuilderImplTest, EmitStatement_Discard) {
auto* expr = Discard();
Func("test_function", {}, ty.void_(), expr,
utils::Vector{
create<ast::StageAttribute>(ast::PipelineStage::kFragment),
});
auto& b = CreateBuilder();
InjectFlowBlock();
b.EmitStatement(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(discard
)");
}
TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) {
Func("my_func", utils::Vector{Param("p", ty.f32())}, ty.void_(), utils::Empty);
auto* stmt = CallStmt(Call("my_func", Mul(2_a, 3_a)));
WrapInFunction(stmt);
auto& b = CreateBuilder();
InjectFlowBlock();
b.EmitStatement(stmt);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:void = call my_func, 6.0f
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
auto i = GlobalVar("i", builtin::AddressSpace::kPrivate, Expr(1_i));
auto* expr = Call(ty.f32(), i);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, i32, read_write> = var private read_write
store %1:ref<private, i32, read_write>, 1i
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
%2:f32 = convert i32, %1:ref<private, i32, read_write>
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
auto* expr = vec3(ty.f32());
GlobalVar("i", builtin::AddressSpace::kPrivate, expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, vec3<f32>, read_write> = var private read_write
store %1:ref<private, vec3<f32>, read_write>, vec3<f32> 0.0f
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
auto i = GlobalVar("i", builtin::AddressSpace::kPrivate, Expr(1_f));
auto* expr = vec3(ty.f32(), 2_f, 3_f, i);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, f32, read_write> = var private read_write
store %1:ref<private, f32, read_write>, 1.0f
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
%2:vec3<f32> = construct 2.0f, 3.0f, %1:ref<private, f32, read_write>
ret
func_end
)");
}
} // namespace
} // namespace tint::ir

115
src/tint/ir/builder_impl_literal_test.cc Normal file
View File

@ -0,0 +1,115 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/ir/test_helper.h"
#include "gmock/gmock.h"
#include "src/tint/ast/case_selector.h"
#include "src/tint/ast/int_literal_expression.h"
#include "src/tint/constant/scalar.h"
namespace tint::ir {
namespace {
using namespace tint::number_suffixes; // NOLINT
using IR_BuilderImplTest = TestHelper;
TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_True) {
auto* expr = Expr(true);
GlobalVar("a", ty.bool_(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<bool>>());
EXPECT_TRUE(val->As<constant::Scalar<bool>>()->ValueAs<bool>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_False) {
auto* expr = Expr(false);
GlobalVar("a", ty.bool_(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<bool>>());
EXPECT_FALSE(val->As<constant::Scalar<bool>>()->ValueAs<bool>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_F32) {
auto* expr = Expr(1.2_f);
GlobalVar("a", ty.f32(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<f32>>());
EXPECT_EQ(1.2_f, val->As<constant::Scalar<f32>>()->ValueAs<f32>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_F16) {
Enable(builtin::Extension::kF16);
auto* expr = Expr(1.2_h);
GlobalVar("a", ty.f16(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<f16>>());
EXPECT_EQ(1.2_h, val->As<constant::Scalar<f16>>()->ValueAs<f32>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_I32) {
auto* expr = Expr(-2_i);
GlobalVar("a", ty.i32(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<i32>>());
EXPECT_EQ(-2_i, val->As<constant::Scalar<i32>>()->ValueAs<f32>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_U32) {
auto* expr = Expr(2_u);
GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<u32>>());
EXPECT_EQ(2_u, val->As<constant::Scalar<u32>>()->ValueAs<f32>());
}
} // namespace
} // namespace tint::ir

47
src/tint/ir/builder_impl_materialize_test.cc Normal file
View File

@ -0,0 +1,47 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/ir/test_helper.h"
#include "gmock/gmock.h"
#include "src/tint/ast/case_selector.h"
#include "src/tint/ast/int_literal_expression.h"
#include "src/tint/constant/scalar.h"
namespace tint::ir {
namespace {
using namespace tint::number_suffixes; // NOLINT
using IR_BuilderImplTest = TestHelper;
TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) {
auto* expr = Return(Call("trunc", 2.5_f));
Func("test_function", {}, ty.f32(), expr, utils::Empty);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function():f32
%fn1 = block
ret 2.0f
func_end
)");
}
} // namespace
} // namespace tint::ir

894
src/tint/ir/builder_impl_test.cc
View File

@ -1494,899 +1494,5 @@ func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_True) {
auto* expr = Expr(true);
GlobalVar("a", ty.bool_(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<bool>>());
EXPECT_TRUE(val->As<constant::Scalar<bool>>()->ValueAs<bool>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_False) {
auto* expr = Expr(false);
GlobalVar("a", ty.bool_(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<bool>>());
EXPECT_FALSE(val->As<constant::Scalar<bool>>()->ValueAs<bool>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_F32) {
auto* expr = Expr(1.2_f);
GlobalVar("a", ty.f32(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<f32>>());
EXPECT_EQ(1.2_f, val->As<constant::Scalar<f32>>()->ValueAs<f32>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_F16) {
Enable(builtin::Extension::kF16);
auto* expr = Expr(1.2_h);
GlobalVar("a", ty.f16(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<f16>>());
EXPECT_EQ(1.2_h, val->As<constant::Scalar<f16>>()->ValueAs<f32>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_I32) {
auto* expr = Expr(-2_i);
GlobalVar("a", ty.i32(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<i32>>());
EXPECT_EQ(-2_i, val->As<constant::Scalar<i32>>()->ValueAs<f32>());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_U32) {
auto* expr = Expr(2_u);
GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate, expr);
auto& b = CreateBuilder();
auto r = b.EmitLiteral(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r.Get()->Is<Constant>());
auto* val = r.Get()->As<Constant>()->value;
EXPECT_TRUE(val->Is<constant::Scalar<u32>>());
EXPECT_EQ(2_u, val->As<constant::Scalar<u32>>()->ValueAs<f32>());
}
TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, u32, read_write> = var private read_write
)");
}
TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
auto* expr = Expr(2_u);
GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate, expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, u32, read_write> = var private read_write
store %1:ref<private, u32, read_write>, 2u
)");
}
TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
auto* a = Var("a", ty.u32(), builtin::AddressSpace::kFunction);
WrapInFunction(a);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
%1:ref<function, u32, read_write> = var function read_write
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
auto* expr = Expr(2_u);
auto* a = Var("a", ty.u32(), builtin::AddressSpace::kFunction, expr);
WrapInFunction(a);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
%1:ref<function, u32, read_write> = var function read_write
store %1:ref<function, u32, read_write>, 2u
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Add) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Add(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = add %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Subtract) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Sub(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = sub %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Multiply) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Mul(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = mul %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Div) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Div(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = div %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Modulo) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Mod(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = mod %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_And) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = And(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = and %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Or) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Or(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = or %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Xor) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Xor(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = xor %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
Func("my_func", utils::Empty, ty.bool_(), utils::Vector{Return(true)});
auto* expr = LogicalAnd(Call("my_func"), false);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():bool
%fn1 = block
ret true
func_end
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1:bool = call my_func
%2:bool = var function read_write
store %2:bool, %1:bool
branch %fn4
%fn4 = if %1:bool [t: %fn5, f: %fn6, m: %fn7]
# true branch
%fn5 = block
store %2:bool, false
branch %fn7
# if merge
%fn7 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
Func("my_func", utils::Empty, ty.bool_(), utils::Vector{Return(true)});
auto* expr = LogicalOr(Call("my_func"), true);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():bool
%fn1 = block
ret true
func_end
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1:bool = call my_func
%2:bool = var function read_write
store %2:bool, %1:bool
branch %fn4
%fn4 = if %1:bool [t: %fn5, f: %fn6, m: %fn7]
# true branch
# false branch
%fn6 = block
store %2:bool, true
branch %fn7
# if merge
%fn7 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Equal) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Equal(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = eq %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_NotEqual) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = NotEqual(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = neq %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThan) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = LessThan(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = lt %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThan) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = GreaterThan(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = gt %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThanEqual) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = LessThanEqual(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = lte %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThanEqual) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = GreaterThanEqual(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = gte %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftLeft) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Shl(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = shiftl %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftRight) {
Func("my_func", utils::Empty, ty.u32(), utils::Vector{Return(0_u)});
auto* expr = Shr(Call("my_func"), 4_u);
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = shiftr %1:u32, 4u
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
Func("my_func", utils::Empty, ty.f32(), utils::Vector{Return(0_f)});
auto* expr = LogicalAnd(LessThan(Call("my_func"), 2_f),
GreaterThan(2.5_f, Div(Call("my_func"), Mul(2.3_f, Call("my_func")))));
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():f32
%fn1 = block
ret 0.0f
func_end
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1:f32 = call my_func
%2:bool = lt %1:f32, 2.0f
%3:bool = var function read_write
store %3:bool, %2:bool
branch %fn4
%fn4 = if %2:bool [t: %fn5, f: %fn6, m: %fn7]
# true branch
%fn5 = block
%4:f32 = call my_func
%5:f32 = call my_func
%6:f32 = mul 2.29999995231628417969f, %5:f32
%7:f32 = div %4:f32, %6:f32
%8:bool = gt 2.5f, %7:f32
store %3:bool, %8:bool
branch %fn7
# if merge
%fn7 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) {
Func("my_func", utils::Vector{Param("p", ty.bool_())}, ty.bool_(), utils::Vector{Return(true)});
auto* expr = Call("my_func", LogicalAnd(LessThan(2.4_f, 2_f),
GreaterThan(2.5_f, Div(10_f, Mul(2.3_f, 9.4_f)))));
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():bool
%fn1 = block
ret true
func_end
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1:bool = call my_func, false
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Bitcast) {
Func("my_func", utils::Empty, ty.f32(), utils::Vector{Return(0_f)});
auto* expr = Bitcast<f32>(Call("my_func"));
WrapInFunction(expr);
auto& b = CreateBuilder();
InjectFlowBlock();
auto r = b.EmitExpression(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
ASSERT_TRUE(r);
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:f32 = call my_func
%2:f32 = bitcast %1:f32
)");
}
TEST_F(IR_BuilderImplTest, EmitStatement_Discard) {
auto* expr = Discard();
Func("test_function", {}, ty.void_(), expr,
utils::Vector{
create<ast::StageAttribute>(ast::PipelineStage::kFragment),
});
auto& b = CreateBuilder();
InjectFlowBlock();
b.EmitStatement(expr);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(discard
)");
}
TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) {
Func("my_func", utils::Vector{Param("p", ty.f32())}, ty.void_(), utils::Empty);
auto* stmt = CallStmt(Call("my_func", Mul(2_a, 3_a)));
WrapInFunction(stmt);
auto& b = CreateBuilder();
InjectFlowBlock();
b.EmitStatement(stmt);
ASSERT_THAT(b.Diagnostics(), testing::IsEmpty());
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1:void = call my_func, 6.0f
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
auto* expr = vec3(ty.f32());
GlobalVar("i", builtin::AddressSpace::kPrivate, expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, vec3<f32>, read_write> = var private read_write
store %1:ref<private, vec3<f32>, read_write>, vec3<f32> 0.0f
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
auto i = GlobalVar("i", builtin::AddressSpace::kPrivate, Expr(1_f));
auto* expr = vec3(ty.f32(), 2_f, 3_f, i);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, f32, read_write> = var private read_write
store %1:ref<private, f32, read_write>, 1.0f
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
%2:vec3<f32> = construct 2.0f, 3.0f, %1:ref<private, f32, read_write>
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
auto i = GlobalVar("i", builtin::AddressSpace::kPrivate, Expr(1_i));
auto* expr = Call(ty.f32(), i);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, i32, read_write> = var private read_write
store %1:ref<private, i32, read_write>, 1i
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
%2:f32 = convert i32, %1:ref<private, i32, read_write>
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) {
auto* expr = Return(Call("trunc", 2.5_f));
Func("test_function", {}, ty.f32(), expr, utils::Empty);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function():f32
%fn1 = block
ret 2.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) {
auto i = GlobalVar("i", builtin::AddressSpace::kPrivate, Expr(1_f));
auto* expr = Call("asin", i);
WrapInFunction(expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, f32, read_write> = var private read_write
store %1:ref<private, f32, read_write>, 1.0f
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
%2:f32 = asin %1:ref<private, f32, read_write>
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_Vertex) {
Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(0_f, 0_f, 0_f, 0_f))},
utils::Vector{Stage(ast::PipelineStage::kVertex)},
utils::Vector{Builtin(builtin::BuiltinValue::kPosition)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@vertex ra: @position]
%fn1 = block
ret vec4<f32> 0.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_Fragment) {
Func("test", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{Stage(ast::PipelineStage::kFragment)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():void [@fragment]
%fn1 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_Compute) {
Func("test", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{Stage(ast::PipelineStage::kCompute), WorkgroupSize(8_i, 4_i, 2_i)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():void [@compute @workgroup_size(8, 4, 2)]
%fn1 = block
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_Return) {
Func("test", utils::Empty, ty.vec3<f32>(), utils::Vector{Return(vec3<f32>(0_f, 0_f, 0_f))},
utils::Empty);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec3<f32>
%fn1 = block
ret vec3<f32> 0.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPosition) {
Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(1_f, 2_f, 3_f, 4_f))},
utils::Vector{Stage(ast::PipelineStage::kVertex)},
utils::Vector{Builtin(builtin::BuiltinValue::kPosition)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@vertex ra: @position]
%fn1 = block
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPositionInvariant) {
Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(1_f, 2_f, 3_f, 4_f))},
utils::Vector{Stage(ast::PipelineStage::kVertex)},
utils::Vector{Builtin(builtin::BuiltinValue::kPosition), Invariant()});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@vertex ra: @position @invariant]
%fn1 = block
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation) {
Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(1_f, 2_f, 3_f, 4_f))},
utils::Vector{Stage(ast::PipelineStage::kFragment)}, utils::Vector{Location(1_i)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@fragment ra: @location(1)]
%fn1 = block
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnFragDepth) {
Func("test", utils::Empty, ty.f32(), utils::Vector{Return(1_f)},
utils::Vector{Stage(ast::PipelineStage::kFragment)},
utils::Vector{Builtin(builtin::BuiltinValue::kFragDepth)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():f32 [@fragment ra: @frag_depth]
%fn1 = block
ret 1.0f
func_end
)");
}
TEST_F(IR_BuilderImplTest, EmitFunction_ReturnSampleMask) {
Func("test", utils::Empty, ty.u32(), utils::Vector{Return(1_u)},
utils::Vector{Stage(ast::PipelineStage::kFragment)},
utils::Vector{Builtin(builtin::BuiltinValue::kSampleMask)});
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():u32 [@fragment ra: @sample_mask]
%fn1 = block
ret 1u
func_end
)");
}
} // namespace
} // namespace tint::ir

99
src/tint/ir/builder_impl_var_test.cc Normal file
View File

@ -0,0 +1,99 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/ir/test_helper.h"
#include "gmock/gmock.h"
#include "src/tint/ast/case_selector.h"
#include "src/tint/ast/int_literal_expression.h"
#include "src/tint/constant/scalar.h"
namespace tint::ir {
namespace {
using namespace tint::number_suffixes; // NOLINT
using IR_BuilderImplTest = TestHelper;
TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, u32, read_write> = var private read_write
)");
}
TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
auto* expr = Expr(2_u);
GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate, expr);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1:ref<private, u32, read_write> = var private read_write
store %1:ref<private, u32, read_write>, 2u
)");
}
TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
auto* a = Var("a", ty.u32(), builtin::AddressSpace::kFunction);
WrapInFunction(a);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
%1:ref<function, u32, read_write> = var function read_write
ret
func_end
)");
}
TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
auto* expr = Expr(2_u);
auto* a = Var("a", ty.u32(), builtin::AddressSpace::kFunction, expr);
WrapInFunction(a);
auto r = Build();
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
%1:ref<function, u32, read_write> = var function read_write
store %1:ref<function, u32, read_write>, 2u
ret
func_end
)");
}
} // namespace
} // namespace tint::ir