From 63716c55cc4f6ea9872bf736a3d2b1e6d2a19d1f Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Thu, 4 May 2023 18:01:26 +0000 Subject: [PATCH] [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 Reviewed-by: James Price Commit-Queue: Dan Sinclair --- builder_impl_builtin_test.cc | 54 ++ builder_impl_function_test.cc | 179 ++++ src/tint/BUILD.gn | 5 + src/tint/CMakeLists.txt | 5 + src/tint/ir/builder_impl_binary_test.cc | 456 ++++++++++ src/tint/ir/builder_impl_call_test.cc | 153 ++++ src/tint/ir/builder_impl_literal_test.cc | 115 +++ src/tint/ir/builder_impl_materialize_test.cc | 47 + src/tint/ir/builder_impl_test.cc | 894 ------------------- src/tint/ir/builder_impl_var_test.cc | 99 ++ 10 files changed, 1113 insertions(+), 894 deletions(-) create mode 100644 builder_impl_builtin_test.cc create mode 100644 builder_impl_function_test.cc create mode 100644 src/tint/ir/builder_impl_binary_test.cc create mode 100644 src/tint/ir/builder_impl_call_test.cc create mode 100644 src/tint/ir/builder_impl_literal_test.cc create mode 100644 src/tint/ir/builder_impl_materialize_test.cc create mode 100644 src/tint/ir/builder_impl_var_test.cc diff --git a/builder_impl_builtin_test.cc b/builder_impl_builtin_test.cc new file mode 100644 index 0000000000..951ec629f6 --- /dev/null +++ b/builder_impl_builtin_test.cc @@ -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 = var private read_write +store %1:ref, 1.0f + + + +%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] + %fn2 = block + %2:f32 = asin %1:ref + ret +func_end + +)"); +} + +} // namespace +} // namespace tint::ir diff --git a/builder_impl_function_test.cc b/builder_impl_function_test.cc new file mode 100644 index 0000000000..d0c6d98824 --- /dev/null +++ b/builder_impl_function_test.cc @@ -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(), utils::Vector{Return(vec4(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 [@vertex ra: @position] + %fn1 = block + ret vec4 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(), utils::Vector{Return(vec3(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 + %fn1 = block + ret vec3 0.0f +func_end + +)"); +} + +TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPosition) { + Func("test", utils::Empty, ty.vec4(), utils::Vector{Return(vec4(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 [@vertex ra: @position] + %fn1 = block + ret vec4 1.0f, 2.0f, 3.0f, 4.0f +func_end + +)"); +} + +TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPositionInvariant) { + Func("test", utils::Empty, ty.vec4(), utils::Vector{Return(vec4(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 [@vertex ra: @position @invariant] + %fn1 = block + ret vec4 1.0f, 2.0f, 3.0f, 4.0f +func_end + +)"); +} + +TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation) { + Func("test", utils::Empty, ty.vec4(), utils::Vector{Return(vec4(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 [@fragment ra: @location(1)] + %fn1 = block + ret vec4 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 diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index c96b913cc4..4bf1bcd2a6 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -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", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index f79de57bd1..2012b6a944 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -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 diff --git a/src/tint/ir/builder_impl_binary_test.cc b/src/tint/ir/builder_impl_binary_test.cc new file mode 100644 index 0000000000..dd8a896eda --- /dev/null +++ b/src/tint/ir/builder_impl_binary_test.cc @@ -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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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()); + 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 diff --git a/src/tint/ir/builder_impl_call_test.cc b/src/tint/ir/builder_impl_call_test.cc new file mode 100644 index 0000000000..456488054f --- /dev/null +++ b/src/tint/ir/builder_impl_call_test.cc @@ -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(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()); + 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::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()); + 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()); + 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 = var private read_write +store %1:ref, 1i + + + +%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] + %fn2 = block + %2:f32 = convert i32, %1:ref + 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, read_write> = var private read_write +store %1:ref, read_write>, vec3 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 = var private read_write +store %1:ref, 1.0f + + + +%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] + %fn2 = block + %2:vec3 = construct 2.0f, 3.0f, %1:ref + ret +func_end + +)"); +} + +} // namespace +} // namespace tint::ir diff --git a/src/tint/ir/builder_impl_literal_test.cc b/src/tint/ir/builder_impl_literal_test.cc new file mode 100644 index 0000000000..090f3f82eb --- /dev/null +++ b/src/tint/ir/builder_impl_literal_test.cc @@ -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()); + auto* val = r.Get()->As()->value; + EXPECT_TRUE(val->Is>()); + EXPECT_TRUE(val->As>()->ValueAs()); +} + +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()); + auto* val = r.Get()->As()->value; + EXPECT_TRUE(val->Is>()); + EXPECT_FALSE(val->As>()->ValueAs()); +} + +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()); + auto* val = r.Get()->As()->value; + EXPECT_TRUE(val->Is>()); + EXPECT_EQ(1.2_f, val->As>()->ValueAs()); +} + +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()); + auto* val = r.Get()->As()->value; + EXPECT_TRUE(val->Is>()); + EXPECT_EQ(1.2_h, val->As>()->ValueAs()); +} + +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()); + auto* val = r.Get()->As()->value; + EXPECT_TRUE(val->Is>()); + EXPECT_EQ(-2_i, val->As>()->ValueAs()); +} + +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()); + auto* val = r.Get()->As()->value; + EXPECT_TRUE(val->Is>()); + EXPECT_EQ(2_u, val->As>()->ValueAs()); +} + +} // namespace +} // namespace tint::ir diff --git a/src/tint/ir/builder_impl_materialize_test.cc b/src/tint/ir/builder_impl_materialize_test.cc new file mode 100644 index 0000000000..4f917f70b7 --- /dev/null +++ b/src/tint/ir/builder_impl_materialize_test.cc @@ -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 diff --git a/src/tint/ir/builder_impl_test.cc b/src/tint/ir/builder_impl_test.cc index ea52891c6e..4a6e713853 100644 --- a/src/tint/ir/builder_impl_test.cc +++ b/src/tint/ir/builder_impl_test.cc @@ -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()); - auto* val = r.Get()->As()->value; - EXPECT_TRUE(val->Is>()); - EXPECT_TRUE(val->As>()->ValueAs()); -} - -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()); - auto* val = r.Get()->As()->value; - EXPECT_TRUE(val->Is>()); - EXPECT_FALSE(val->As>()->ValueAs()); -} - -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()); - auto* val = r.Get()->As()->value; - EXPECT_TRUE(val->Is>()); - EXPECT_EQ(1.2_f, val->As>()->ValueAs()); -} - -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()); - auto* val = r.Get()->As()->value; - EXPECT_TRUE(val->Is>()); - EXPECT_EQ(1.2_h, val->As>()->ValueAs()); -} - -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()); - auto* val = r.Get()->As()->value; - EXPECT_TRUE(val->Is>()); - EXPECT_EQ(-2_i, val->As>()->ValueAs()); -} - -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()); - auto* val = r.Get()->As()->value; - EXPECT_TRUE(val->Is>()); - EXPECT_EQ(2_u, val->As>()->ValueAs()); -} - -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 = 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 = var private read_write -store %1:ref, 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 = 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 = var function read_write - store %1:ref, 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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()); - 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(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()); - 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::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()); - 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()); - 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, read_write> = var private read_write -store %1:ref, read_write>, vec3 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 = var private read_write -store %1:ref, 1.0f - - - -%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] - %fn2 = block - %2:vec3 = construct 2.0f, 3.0f, %1:ref - 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 = var private read_write -store %1:ref, 1i - - - -%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] - %fn2 = block - %2:f32 = convert i32, %1:ref - 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 = var private read_write -store %1:ref, 1.0f - - - -%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] - %fn2 = block - %2:f32 = asin %1:ref - ret -func_end - -)"); -} - -TEST_F(IR_BuilderImplTest, EmitFunction_Vertex) { - Func("test", utils::Empty, ty.vec4(), utils::Vector{Return(vec4(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 [@vertex ra: @position] - %fn1 = block - ret vec4 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(), utils::Vector{Return(vec3(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 - %fn1 = block - ret vec3 0.0f -func_end - -)"); -} - -TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPosition) { - Func("test", utils::Empty, ty.vec4(), utils::Vector{Return(vec4(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 [@vertex ra: @position] - %fn1 = block - ret vec4 1.0f, 2.0f, 3.0f, 4.0f -func_end - -)"); -} - -TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPositionInvariant) { - Func("test", utils::Empty, ty.vec4(), utils::Vector{Return(vec4(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 [@vertex ra: @position @invariant] - %fn1 = block - ret vec4 1.0f, 2.0f, 3.0f, 4.0f -func_end - -)"); -} - -TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation) { - Func("test", utils::Empty, ty.vec4(), utils::Vector{Return(vec4(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 [@fragment ra: @location(1)] - %fn1 = block - ret vec4 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 diff --git a/src/tint/ir/builder_impl_var_test.cc b/src/tint/ir/builder_impl_var_test.cc new file mode 100644 index 0000000000..485684a4b3 --- /dev/null +++ b/src/tint/ir/builder_impl_var_test.cc @@ -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 = 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 = var private read_write +store %1:ref, 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 = 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 = var function read_write + store %1:ref, 2u + ret +func_end + +)"); +} +} // namespace +} // namespace tint::ir