| // Copyright 2023 The Dawn & Tint Authors |
| // |
| // Redistribution and use in source and binary forms, with or without |
| // modification, are permitted provided that the following conditions are met: |
| // |
| // 1. Redistributions of source code must retain the above copyright notice, this |
| // list of conditions and the following disclaimer. |
| // |
| // 2. Redistributions in binary form must reproduce the above copyright notice, |
| // this list of conditions and the following disclaimer in the documentation |
| // and/or other materials provided with the distribution. |
| // |
| // 3. Neither the name of the copyright holder nor the names of its |
| // contributors may be used to endorse or promote products derived from |
| // this software without specific prior written permission. |
| // |
| // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" |
| // AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE |
| // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE |
| // DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE |
| // FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL |
| // DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR |
| // SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER |
| // CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, |
| // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE |
| // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
| |
| #include "gmock/gmock.h" |
| #include "src/tint/lang/core/builtin_fn.h" |
| #include "src/tint/lang/core/constant/scalar.h" |
| #include "src/tint/lang/core/fluent_types.h" |
| #include "src/tint/lang/core/ir/block.h" |
| #include "src/tint/lang/core/ir/if.h" |
| #include "src/tint/lang/core/ir/loop.h" |
| #include "src/tint/lang/core/ir/multi_in_block.h" |
| #include "src/tint/lang/core/ir/switch.h" |
| #include "src/tint/lang/wgsl/ast/case_selector.h" |
| #include "src/tint/lang/wgsl/ast/int_literal_expression.h" |
| #include "src/tint/lang/wgsl/helpers/ir_program_test.h" |
| |
| using namespace tint::core::fluent_types; // NOLINT |
| |
| namespace tint::wgsl::reader { |
| namespace { |
| |
| /// Looks for the instruction with the given type T. |
| /// If no instruction is found, then nullptr is returned. |
| /// If multiple instructions are found with the type T, then an error is raised and the first is |
| /// returned. |
| template <typename T> |
| T* FindSingleInstruction(core::ir::Module& mod) { |
| T* found = nullptr; |
| size_t count = 0; |
| for (auto* node : mod.Instructions()) { |
| if (auto* as = node->As<T>()) { |
| count++; |
| if (!found) { |
| found = as; |
| } |
| } |
| } |
| if (count > 1) { |
| ADD_FAILURE() << "FindSingleInstruction() found " << count << " nodes of type " |
| << tint::TypeInfo::Of<T>().name; |
| } |
| return found; |
| } |
| |
| using namespace tint::core::number_suffixes; // NOLINT |
| |
| using IR_FromProgramTest = helpers::IRProgramTest; |
| |
| TEST_F(IR_FromProgramTest, Func) { |
| Func("f", tint::Empty, ty.void_(), tint::Empty); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| ASSERT_EQ(1u, m->functions.Length()); |
| |
| core::ir::Function* f = m->functions[0]; |
| ASSERT_NE(f->Block(), nullptr); |
| |
| EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kUndefined); |
| |
| EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():void { |
| $B1: { |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Func_WithParam) { |
| Func("f", Vector{Param("a", ty.u32())}, ty.u32(), Vector{Return("a")}); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| ASSERT_EQ(1u, m->functions.Length()); |
| |
| core::ir::Function* f = m->functions[0]; |
| ASSERT_NE(f->Block(), nullptr); |
| |
| EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kUndefined); |
| |
| EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func(%a:u32):u32 { |
| $B1: { |
| ret %a |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Func_WithMultipleParam) { |
| Func("f", Vector{Param("a", ty.u32()), Param("b", ty.i32()), Param("c", ty.bool_())}, |
| ty.void_(), tint::Empty); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| ASSERT_EQ(1u, m->functions.Length()); |
| |
| core::ir::Function* f = m->functions[0]; |
| ASSERT_NE(f->Block(), nullptr); |
| |
| EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kUndefined); |
| |
| EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func(%a:u32, %b:i32, %c:bool):void { |
| $B1: { |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, EntryPoint) { |
| Func("f", tint::Empty, ty.void_(), tint::Empty, Vector{Stage(ast::PipelineStage::kFragment)}); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kFragment); |
| } |
| |
| TEST_F(IR_FromProgramTest, IfStatement) { |
| auto* ast_if = If(true, Block(), Else(Block())); |
| WrapInFunction(ast_if); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| exit_if # if_1 |
| } |
| $B3: { # false |
| exit_if # if_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, IfStatement_TrueReturns) { |
| auto* ast_if = If(true, Block(Return())); |
| WrapInFunction(ast_if); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| if true [t: $B2] { # if_1 |
| $B2: { # true |
| ret |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, IfStatement_FalseReturns) { |
| auto* ast_if = If(true, Block(), Else(Block(Return()))); |
| WrapInFunction(ast_if); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| exit_if # if_1 |
| } |
| $B3: { # false |
| ret |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, IfStatement_BothReturn) { |
| auto* ast_if = If(true, Block(Return()), Else(Block(Return()))); |
| WrapInFunction(ast_if); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| if true [t: $B2, f: $B3] { # if_1 |
| $B2: { # true |
| ret |
| } |
| $B3: { # false |
| ret |
| } |
| } |
| unreachable |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, IfStatement_JumpChainToMerge) { |
| auto* ast_loop = Loop(Block(Break())); |
| auto* ast_if = If(true, Block(ast_loop)); |
| WrapInFunction(ast_if); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| if true [t: $B2] { # if_1 |
| $B2: { # true |
| loop [b: $B3] { # loop_1 |
| $B3: { # body |
| exit_loop # loop_1 |
| } |
| } |
| exit_if # if_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Loop_WithBreak) { |
| auto* ast_loop = Loop(Block(Break())); |
| WrapInFunction(ast_loop); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2] { # loop_1 |
| $B2: { # body |
| exit_loop # loop_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Loop_WithContinue) { |
| auto* ast_if = If(true, Block(Break())); |
| auto* ast_loop = Loop(Block(ast_if, Continue())); |
| WrapInFunction(ast_loop); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4] { # if_1 |
| $B4: { # true |
| exit_loop # loop_1 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Loop_WithContinuing_BreakIf) { |
| auto* ast_break_if = BreakIf(true); |
| auto* ast_loop = Loop(Block(), Block(ast_break_if)); |
| WrapInFunction(ast_loop); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| break_if true # -> [t: exit_loop loop_1, f: $B2] |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Loop_Continuing_Body_Scope) { |
| auto* a = Decl(Let("a", Expr(true))); |
| auto* ast_break_if = BreakIf("a"); |
| auto* ast_loop = Loop(Block(a), Block(ast_break_if)); |
| WrapInFunction(ast_loop); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| %a:bool = let true |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| break_if %a # -> [t: exit_loop loop_1, f: $B2] |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Loop_WithReturn) { |
| auto* ast_if = If(true, Block(Return())); |
| auto* ast_loop = Loop(Block(ast_if, Continue())); |
| WrapInFunction(ast_loop); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4] { # if_1 |
| $B4: { # true |
| ret |
| } |
| } |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| unreachable |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Loop_WithOnlyReturn) { |
| auto* ast_loop = Loop(Block(Return(), Continue())); |
| WrapInFunction(ast_loop, If(true, Block(Return()))); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2] { # loop_1 |
| $B2: { # body |
| ret |
| } |
| } |
| unreachable |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Loop_WithOnlyReturn_ContinuingBreakIf) { |
| // Note, even though there is code in the loop merge (specifically, the |
| // `ast_if` below), it doesn't get emitted as there is no way to reach the |
| // loop merge due to the loop itself doing a `return`. This is why the |
| // loop merge gets marked as Dead and the `ast_if` doesn't appear. |
| // |
| // Similar, the continuing block goes away as there is no way to get there, so it's treated |
| // as dead code and dropped. |
| auto* ast_break_if = BreakIf(true); |
| auto* ast_loop = Loop(Block(Return()), Block(ast_break_if)); |
| auto* ast_if = If(true, Block(Return())); |
| WrapInFunction(Block(ast_loop, ast_if)); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2] { # loop_1 |
| $B2: { # body |
| ret |
| } |
| } |
| if true [t: $B3] { # if_1 |
| $B3: { # true |
| ret |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Loop_WithIf_BothBranchesBreak) { |
| auto* ast_if = If(true, Block(Break()), Else(Block(Break()))); |
| auto* ast_loop = Loop(Block(ast_if, Continue())); |
| WrapInFunction(ast_loop); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2] { # loop_1 |
| $B2: { # body |
| if true [t: $B3, f: $B4] { # if_1 |
| $B3: { # true |
| exit_loop # loop_1 |
| } |
| $B4: { # false |
| exit_loop # loop_1 |
| } |
| } |
| unreachable |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Loop_Nested) { |
| auto* ast_if_a = If(true, Block(Break())); |
| auto* ast_if_b = If(true, Block(Continue())); |
| auto* ast_if_c = BreakIf(true); |
| auto* ast_if_d = If(true, Block(Break())); |
| |
| auto* ast_loop_d = Loop(Block(), Block(ast_if_c)); |
| auto* ast_loop_c = Loop(Block(Break())); |
| |
| auto* ast_loop_b = Loop(Block(ast_if_a, ast_if_b), Block(ast_loop_c, ast_loop_d)); |
| auto* ast_loop_a = Loop(Block(ast_loop_b, ast_if_d)); |
| |
| WrapInFunction(ast_loop_a); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| EXPECT_EQ(Disassemble(m.Get()).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| loop [b: $B4, c: $B5] { # loop_2 |
| $B4: { # body |
| if true [t: $B6] { # if_1 |
| $B6: { # true |
| exit_loop # loop_2 |
| } |
| } |
| if true [t: $B7] { # if_2 |
| $B7: { # true |
| continue # -> $B5 |
| } |
| } |
| continue # -> $B5 |
| } |
| $B5: { # continuing |
| loop [b: $B8] { # loop_3 |
| $B8: { # body |
| exit_loop # loop_3 |
| } |
| } |
| loop [b: $B9, c: $B10] { # loop_4 |
| $B9: { # body |
| continue # -> $B10 |
| } |
| $B10: { # continuing |
| break_if true # -> [t: exit_loop loop_4, f: $B9] |
| } |
| } |
| next_iteration # -> $B4 |
| } |
| } |
| if true [t: $B11] { # if_3 |
| $B11: { # true |
| exit_loop # loop_1 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, While) { |
| auto* ast_while = While(false, Block()); |
| WrapInFunction(ast_while); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if false [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| continue # -> $B3 |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, While_Return) { |
| auto* ast_while = While(true, Block(Return())); |
| WrapInFunction(ast_while); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2, c: $B3] { # loop_1 |
| $B2: { # body |
| if true [t: $B4, f: $B5] { # if_1 |
| $B4: { # true |
| exit_if # if_1 |
| } |
| $B5: { # false |
| exit_loop # loop_1 |
| } |
| } |
| ret |
| } |
| $B3: { # continuing |
| next_iteration # -> $B2 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, For) { |
| auto* ast_for = For(Decl(Var("i", ty.i32())), LessThan("i", 10_a), Increment("i"), Block()); |
| WrapInFunction(ast_for); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(2u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [i: $B2, b: $B3, c: $B4] { # loop_1 |
| $B2: { # initializer |
| %i:ptr<function, i32, read_write> = var |
| next_iteration # -> $B3 |
| } |
| $B3: { # body |
| %3:i32 = load %i |
| %4:bool = lt %3, 10i |
| if %4 [t: $B5, f: $B6] { # if_1 |
| $B5: { # true |
| exit_if # if_1 |
| } |
| $B6: { # false |
| exit_loop # loop_1 |
| } |
| } |
| continue # -> $B4 |
| } |
| $B4: { # continuing |
| %5:i32 = load %i |
| %6:i32 = add %5, 1i |
| store %i, %6 |
| next_iteration # -> $B3 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, For_Init_NoCondOrContinuing) { |
| auto* ast_for = For(Decl(Var("i", ty.i32())), nullptr, nullptr, Block(Break())); |
| WrapInFunction(ast_for); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [i: $B2, b: $B3] { # loop_1 |
| $B2: { # initializer |
| %i:ptr<function, i32, read_write> = var |
| next_iteration # -> $B3 |
| } |
| $B3: { # body |
| exit_loop # loop_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, For_NoInitCondOrContinuing) { |
| auto* ast_for = For(nullptr, nullptr, nullptr, Block(Break())); |
| WrapInFunction(ast_for); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* loop = FindSingleInstruction<core::ir::Loop>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length()); |
| EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| loop [b: $B2] { # loop_1 |
| $B2: { # body |
| exit_loop # loop_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Switch) { |
| auto* ast_switch = |
| Switch(1_i, Vector{Case(Vector{CaseSelector(0_i)}, Block()), |
| Case(Vector{CaseSelector(1_i)}, Block()), DefaultCase(Block())}); |
| |
| WrapInFunction(ast_switch); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* swtch = FindSingleInstruction<core::ir::Switch>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| auto& cases = swtch->Cases(); |
| ASSERT_EQ(3u, cases.Length()); |
| |
| ASSERT_EQ(1u, cases[0].selectors.Length()); |
| ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<core::constant::Scalar<i32>>()); |
| EXPECT_EQ(0_i, |
| cases[0].selectors[0].val->Value()->As<core::constant::Scalar<i32>>()->ValueOf()); |
| |
| ASSERT_EQ(1u, cases[1].selectors.Length()); |
| ASSERT_TRUE(cases[1].selectors[0].val->Value()->Is<core::constant::Scalar<i32>>()); |
| EXPECT_EQ(1_i, |
| cases[1].selectors[0].val->Value()->As<core::constant::Scalar<i32>>()->ValueOf()); |
| |
| ASSERT_EQ(1u, cases[2].selectors.Length()); |
| EXPECT_TRUE(cases[2].selectors[0].IsDefault()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| switch 1i [c: (0i, $B2), c: (1i, $B3), c: (default, $B4)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| exit_switch # switch_1 |
| } |
| $B4: { # case |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Switch_MultiSelector) { |
| auto* ast_switch = Switch( |
| 1_i, |
| Vector{Case(Vector{CaseSelector(0_i), CaseSelector(1_i), DefaultCaseSelector()}, Block())}); |
| |
| WrapInFunction(ast_switch); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* swtch = FindSingleInstruction<core::ir::Switch>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| auto& cases = swtch->Cases(); |
| ASSERT_EQ(1u, cases.Length()); |
| ASSERT_EQ(3u, cases[0].selectors.Length()); |
| ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<core::constant::Scalar<i32>>()); |
| EXPECT_EQ(0_i, |
| cases[0].selectors[0].val->Value()->As<core::constant::Scalar<i32>>()->ValueOf()); |
| |
| ASSERT_TRUE(cases[0].selectors[1].val->Value()->Is<core::constant::Scalar<i32>>()); |
| EXPECT_EQ(1_i, |
| cases[0].selectors[1].val->Value()->As<core::constant::Scalar<i32>>()->ValueOf()); |
| |
| EXPECT_TRUE(cases[0].selectors[2].IsDefault()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| switch 1i [c: (0i 1i default, $B2)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Switch_OnlyDefault) { |
| auto* ast_switch = Switch(1_i, Vector{DefaultCase(Block())}); |
| WrapInFunction(ast_switch); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* swtch = FindSingleInstruction<core::ir::Switch>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| auto& cases = swtch->Cases(); |
| ASSERT_EQ(1u, cases.Length()); |
| ASSERT_EQ(1u, cases[0].selectors.Length()); |
| EXPECT_TRUE(cases[0].selectors[0].IsDefault()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| switch 1i [c: (default, $B2)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Switch_WithBreak) { |
| auto* ast_switch = Switch( |
| 1_i, Vector{Case(Vector{CaseSelector(0_i)}, Block(Break(), If(true, Block(Return())))), |
| DefaultCase(Block())}); |
| WrapInFunction(ast_switch); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| auto* swtch = FindSingleInstruction<core::ir::Switch>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| auto& cases = swtch->Cases(); |
| ASSERT_EQ(2u, cases.Length()); |
| ASSERT_EQ(1u, cases[0].selectors.Length()); |
| ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<core::constant::Scalar<i32>>()); |
| EXPECT_EQ(0_i, |
| cases[0].selectors[0].val->Value()->As<core::constant::Scalar<i32>>()->ValueOf()); |
| |
| ASSERT_EQ(1u, cases[1].selectors.Length()); |
| EXPECT_TRUE(cases[1].selectors[0].IsDefault()); |
| |
| // This is 1 because the if is dead-code eliminated and the return doesn't happen. |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| switch 1i [c: (0i, $B2), c: (default, $B3)] { # switch_1 |
| $B2: { # case |
| exit_switch # switch_1 |
| } |
| $B3: { # case |
| exit_switch # switch_1 |
| } |
| } |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Switch_AllReturn) { |
| auto* ast_switch = Switch(1_i, Vector{Case(Vector{CaseSelector(0_i)}, Block(Return())), |
| DefaultCase(Block(Return()))}); |
| auto* ast_if = If(true, Block(Return())); |
| WrapInFunction(ast_switch, ast_if); |
| |
| auto res = Build(); |
| ASSERT_EQ(res, Success); |
| |
| auto m = res.Move(); |
| |
| auto* swtch = FindSingleInstruction<core::ir::Switch>(m); |
| |
| ASSERT_EQ(1u, m.functions.Length()); |
| |
| auto& cases = swtch->Cases(); |
| ASSERT_EQ(2u, cases.Length()); |
| ASSERT_EQ(1u, cases[0].selectors.Length()); |
| ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<core::constant::Scalar<i32>>()); |
| EXPECT_EQ(0_i, |
| cases[0].selectors[0].val->Value()->As<core::constant::Scalar<i32>>()->ValueOf()); |
| |
| ASSERT_EQ(1u, cases[1].selectors.Length()); |
| EXPECT_TRUE(cases[1].selectors[0].IsDefault()); |
| |
| EXPECT_EQ(Disassemble(m).Plain(), |
| R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B1: { |
| switch 1i [c: (0i, $B2), c: (default, $B3)] { # switch_1 |
| $B2: { # case |
| ret |
| } |
| $B3: { # case |
| ret |
| } |
| } |
| unreachable |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Emit_Phony) { |
| Func("b", tint::Empty, ty.i32(), Return(1_i)); |
| WrapInFunction(Ignore(Call("b"))); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| EXPECT_EQ(Disassemble(m.Get()).Plain(), |
| R"(%b = func():i32 { |
| $B1: { |
| ret 1i |
| } |
| } |
| %test_function = @compute @workgroup_size(1, 1, 1) func():void { |
| $B2: { |
| %3:i32 = call %b |
| ret |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Func_WithParam_WithAttribute_Invariant) { |
| Func("f", |
| Vector{Param("a", ty.vec4<f32>(), |
| Vector{Invariant(), Builtin(core::BuiltinValue::kPosition)})}, |
| ty.vec4<f32>(), Vector{Return("a")}, Vector{Stage(ast::PipelineStage::kFragment)}, |
| Vector{Location(1_i)}); |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| EXPECT_EQ( |
| Disassemble(m.Get()).Plain(), |
| R"(%f = @fragment func(%a:vec4<f32> [@invariant, @position]):vec4<f32> [@location(1)] { |
| $B1: { |
| ret %a |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Func_WithParam_WithAttribute_Location) { |
| Func("f", Vector{Param("a", ty.f32(), Vector{Location(2_i)})}, ty.f32(), Vector{Return("a")}, |
| Vector{Stage(ast::PipelineStage::kFragment)}, Vector{Location(1_i)}); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| EXPECT_EQ(Disassemble(m.Get()).Plain(), |
| R"(%f = @fragment func(%a:f32 [@location(2)]):f32 [@location(1)] { |
| $B1: { |
| ret %a |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Func_WithParam_WithAttribute_Location_WithInterpolation_LinearCentroid) { |
| Func("f", |
| Vector{Param("a", ty.f32(), |
| Vector{Location(2_i), Interpolate(core::InterpolationType::kLinear, |
| core::InterpolationSampling::kCentroid)})}, |
| ty.f32(), Vector{Return("a")}, Vector{Stage(ast::PipelineStage::kFragment)}, |
| Vector{Location(1_i)}); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| EXPECT_EQ( |
| Disassemble(m.Get()).Plain(), |
| R"(%f = @fragment func(%a:f32 [@location(2), @interpolate(linear, centroid)]):f32 [@location(1)] { |
| $B1: { |
| ret %a |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Func_WithParam_WithAttribute_Location_WithInterpolation_Flat) { |
| Func("f", |
| Vector{Param("a", ty.f32(), |
| Vector{Location(2_i), Interpolate(core::InterpolationType::kFlat)})}, |
| ty.f32(), Vector{Return("a")}, Vector{Stage(ast::PipelineStage::kFragment)}, |
| Vector{Location(1_i)}); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| EXPECT_EQ(Disassemble(m.Get()).Plain(), |
| R"(%f = @fragment func(%a:f32 [@location(2), @interpolate(flat)]):f32 [@location(1)] { |
| $B1: { |
| ret %a |
| } |
| } |
| )"); |
| } |
| |
| TEST_F(IR_FromProgramTest, Requires) { |
| Require(wgsl::LanguageFeature::kReadonlyAndReadwriteStorageTextures); |
| Func("f", tint::Empty, ty.void_(), tint::Empty); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| ASSERT_EQ(1u, m->functions.Length()); |
| |
| core::ir::Function* f = m->functions[0]; |
| ASSERT_NE(f->Block(), nullptr); |
| |
| EXPECT_EQ(m->functions[0]->Stage(), core::ir::Function::PipelineStage::kUndefined); |
| |
| EXPECT_EQ(Disassemble(m.Get()).Plain(), R"(%f = func():void { |
| $B1: { |
| ret |
| } |
| } |
| )"); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| // Bugs |
| //////////////////////////////////////////////////////////////////////////////// |
| TEST_F(IR_FromProgramTest, BugChromium324466107) { |
| Func("f", Empty, ty.void_(), |
| Vector{ |
| // Abstract type on the RHS - cannot be emitted. |
| Assign(Phony(), Call(core::BuiltinFn::kFrexp, Call(ty.vec2<Infer>(), 2.0_a))), |
| }); |
| |
| auto m = Build(); |
| ASSERT_EQ(m, Success); |
| |
| EXPECT_EQ(Disassemble(m.Get()).Plain(), |
| R"(%f = func():void { |
| $B1: { |
| ret |
| } |
| } |
| )"); |
| } |
| |
| } // namespace |
| } // namespace tint::wgsl::reader |