From 69bb5dd816781b7490c9a3cc6a5de1f186bc9f2f Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Wed, 3 May 2023 21:31:51 +0000 Subject: [PATCH] [ir] Add function return information. This Cl adds information into the IR on the function return type and any associated attributes. Bug: tint:1915 Change-Id: I74cbf2613b4ae575e33a61d04d30b515df6ba796 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/131300 Reviewed-by: James Price Commit-Queue: Dan Sinclair Kokoro: Kokoro --- src/tint/ir/builder_impl.cc | 70 +++++++++--- src/tint/ir/builder_impl_test.cc | 178 ++++++++++++++++++++++++------- src/tint/ir/disassembler.cc | 17 ++- src/tint/ir/function.cc | 18 ++++ src/tint/ir/function.h | 25 +++++ 5 files changed, 254 insertions(+), 54 deletions(-) diff --git a/src/tint/ir/builder_impl.cc b/src/tint/ir/builder_impl.cc index eedf34490a..6ee8170224 100644 --- a/src/tint/ir/builder_impl.cc +++ b/src/tint/ir/builder_impl.cc @@ -39,6 +39,7 @@ #include "src/tint/ast/identifier_expression.h" #include "src/tint/ast/if_statement.h" #include "src/tint/ast/int_literal_expression.h" +#include "src/tint/ast/invariant_attribute.h" #include "src/tint/ast/let.h" #include "src/tint/ast/literal_expression.h" #include "src/tint/ast/loop_statement.h" @@ -211,6 +212,7 @@ void BuilderImpl::EmitFunction(const ast::Function* ast_func) { ast_to_flow_[ast_func] = ir_func; + const auto* sem = program_->Sem().Get(ast_func); if (ast_func->IsEntryPoint()) { builder.ir.entry_points.Push(ir_func); @@ -224,7 +226,6 @@ void BuilderImpl::EmitFunction(const ast::Function* ast_func) { case ast::PipelineStage::kCompute: { ir_func->pipeline_stage = Function::PipelineStage::kCompute; - const auto* sem = program_->Sem().Get(ast_func); auto wg_size = sem->WorkgroupSize(); uint32_t x = wg_size[0].value(); @@ -246,7 +247,49 @@ void BuilderImpl::EmitFunction(const ast::Function* ast_func) { return; } } + + for (auto* attr : ast_func->return_type_attributes) { + tint::Switch( + attr, // + [&](const ast::LocationAttribute*) { + ir_func->return_attributes.Push(Function::ReturnAttribute::kLocation); + }, + [&](const ast::InvariantAttribute*) { + ir_func->return_attributes.Push(Function::ReturnAttribute::kInvariant); + }, + [&](const ast::BuiltinAttribute* b) { + if (auto* ident_sem = + program_->Sem() + .Get(b) + ->As>()) { + switch (ident_sem->Value()) { + case builtin::BuiltinValue::kPosition: + ir_func->return_attributes.Push( + Function::ReturnAttribute::kPosition); + break; + case builtin::BuiltinValue::kFragDepth: + ir_func->return_attributes.Push( + Function::ReturnAttribute::kFragDepth); + break; + case builtin::BuiltinValue::kSampleMask: + ir_func->return_attributes.Push( + Function::ReturnAttribute::kSampleMask); + break; + default: + TINT_ICE(IR, diagnostics_) + << "Unknown builtin value in return attributes " + << ident_sem->Value(); + return; + } + } else { + TINT_ICE(IR, diagnostics_) << "Builtin attribute sem invalid"; + return; + } + }); + } } + ir_func->return_type = sem->ReturnType()->Clone(clone_ctx_.type_ctx); + ir_func->return_location = sem->ReturnLocation(); { FlowStackScope scope(this, ir_func); @@ -400,9 +443,9 @@ void BuilderImpl::EmitBlock(const ast::BlockStatement* block) { scopes_.Push(); TINT_DEFER(scopes_.Pop()); - // Note, this doesn't need to emit a Block as the current block flow node should be - // sufficient as the blocks all get flattened. Each flow control node will inject the basic - // blocks it requires. + // Note, this doesn't need to emit a Block as the current block flow node should be sufficient + // as the blocks all get flattened. Each flow control node will inject the basic blocks it + // requires. EmitStatements(block->statements); } @@ -439,9 +482,8 @@ void BuilderImpl::EmitIf(const ast::IfStatement* stmt) { } current_flow_block = nullptr; - // If both branches went somewhere, then they both returned, continued or broke. So, - // there is no need for the if merge-block and there is nothing to branch to the merge - // block anyway. + // If both branches went somewhere, then they both returned, continued or broke. So, there is no + // need for the if merge-block and there is nothing to branch to the merge block anyway. if (IsConnected(if_node->merge.target)) { current_flow_block = if_node->merge.target->As(); } @@ -472,8 +514,8 @@ void BuilderImpl::EmitLoop(const ast::LoopStatement* stmt) { BranchToIfNeeded(loop_node->start.target); } - // The loop merge can get disconnected if the loop returns directly, or the continuing - // target branches, eventually, to the merge, but nothing branched to the continuing target. + // The loop merge can get disconnected if the loop returns directly, or the continuing target + // branches, eventually, to the merge, but nothing branched to the continuing target. current_flow_block = loop_node->merge.target->As(); if (!IsConnected(loop_node->merge.target)) { current_flow_block = nullptr; @@ -661,9 +703,9 @@ void BuilderImpl::EmitContinue(const ast::ContinueStatement*) { } // Discard is being treated as an instruction. The semantics in WGSL is demote_to_helper, so the -// code has to continue as before it just predicates writes. If WGSL grows some kind of -// terminating discard that would probably make sense as a FlowNode but would then require -// figuring out the multi-level exit that is triggered. +// code has to continue as before it just predicates writes. If WGSL grows some kind of terminating +// discard that would probably make sense as a FlowNode but would then require figuring out the +// multi-level exit that is triggered. void BuilderImpl::EmitDiscard(const ast::DiscardStatement*) { auto* inst = builder.Discard(); current_flow_block->instructions.Push(inst); @@ -783,8 +825,8 @@ void BuilderImpl::EmitVariable(const ast::Variable* var) { // should never be used. // // TODO(dsinclair): Probably want to store the const variable somewhere and then in - // identifier expression log an error if we ever see a const identifier. Add this - // when identifiers and variables are supported. + // identifier expression log an error if we ever see a const identifier. Add this when + // identifiers and variables are supported. }, [&](Default) { add_error(var->source, "unknown variable: " + std::string(var->TypeInfo().name)); diff --git a/src/tint/ir/builder_impl_test.cc b/src/tint/ir/builder_impl_test.cc index 431423f554..9fc8edf49f 100644 --- a/src/tint/ir/builder_impl_test.cc +++ b/src/tint/ir/builder_impl_test.cc @@ -42,7 +42,7 @@ TEST_F(IR_BuilderImplTest, Func) { EXPECT_EQ(1u, f->start_target->inbound_branches.Length()); EXPECT_EQ(1u, f->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func f + EXPECT_EQ(Disassemble(m), R"(%fn0 = func f(void) %fn1 = block ret func_end @@ -88,7 +88,7 @@ TEST_F(IR_BuilderImplTest, IfStatement) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -136,7 +136,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(2u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -183,7 +183,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(2u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -230,7 +230,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(2u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -273,7 +273,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) { ASSERT_NE(loop_flow->continuing.target, nullptr); ASSERT_NE(loop_flow->merge.target, nullptr); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -330,7 +330,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -388,7 +388,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -463,7 +463,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -538,7 +538,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -595,7 +595,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -657,7 +657,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) { // This is 1 because only the loop branch happens. The subsequent if return is dead code. EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -711,7 +711,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -857,7 +857,7 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) { EXPECT_EQ(1u, func->start_target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -1001,7 +1001,7 @@ TEST_F(IR_BuilderImplTest, While) { EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length()); EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -1071,7 +1071,7 @@ TEST_F(IR_BuilderImplTest, While_Return) { EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length()); EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -1178,7 +1178,7 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) { EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -1237,7 +1237,7 @@ TEST_F(IR_BuilderImplTest, Switch) { EXPECT_EQ(3u, flow->merge.target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -1301,7 +1301,7 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) { EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -1345,7 +1345,7 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) { EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length()); EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -1398,7 +1398,7 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) { // This is 1 because the if is dead-code eliminated and the return doesn't happen. EXPECT_EQ(1u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -1457,7 +1457,7 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) { EXPECT_EQ(0u, flow->merge.target->inbound_branches.Length()); EXPECT_EQ(2u, func->end_target->inbound_branches.Length()); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn1 = block branch %fn2 @@ -1596,7 +1596,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + 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 @@ -1614,7 +1614,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)] + 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 @@ -1777,12 +1777,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func + EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(bool) %fn1 = block ret true func_end -%fn2 = func test_function [@compute @workgroup_size(1, 1, 1)] +%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn3 = block %1(bool) = call my_func %2(bool) = var function read_write @@ -1812,12 +1812,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func + EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(bool) %fn1 = block ret true func_end -%fn2 = func test_function [@compute @workgroup_size(1, 1, 1)] +%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn3 = block %1(bool) = call my_func %2(bool) = var function read_write @@ -1994,12 +1994,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func + EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(f32) %fn1 = block ret 0.0f func_end -%fn2 = func test_function [@compute @workgroup_size(1, 1, 1)] +%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 @@ -2036,12 +2036,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func + EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(bool) %fn1 = block ret true func_end -%fn2 = func test_function [@compute @workgroup_size(1, 1, 1)] +%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn3 = block %1(bool) = call my_func, false ret @@ -2137,7 +2137,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) { store %1(ref), 1.0f ret -%fn1 = func test_function [@compute @workgroup_size(1, 1, 1)] +%fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn2 = block %2(vec3) = construct 2.0f, 3.0f, %1(ref) ret @@ -2161,7 +2161,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) { store %1(ref), 1i ret -%fn1 = func test_function [@compute @workgroup_size(1, 1, 1)] +%fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn2 = block %2(f32) = convert i32, %1(ref) ret @@ -2179,7 +2179,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(f32) %fn1 = block ret 2.0f func_end @@ -2201,7 +2201,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) { store %1(ref), 1.0f ret -%fn1 = func test_function [@compute @workgroup_size(1, 1, 1)] +%fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn2 = block %2(f32) = asin %1(ref) ret @@ -2219,7 +2219,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Vertex) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test [@vertex] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec4) [@vertex ra: @position] %fn1 = block ret vec4 0.0f func_end @@ -2235,7 +2235,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Fragment) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test [@fragment] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(void) [@fragment] %fn1 = block ret func_end @@ -2251,7 +2251,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Compute) { ASSERT_TRUE(r) << Error(); auto m = r.Move(); - EXPECT_EQ(Disassemble(m), R"(%fn0 = func test [@compute @workgroup_size(8, 4, 2)] + EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(void) [@compute @workgroup_size(8, 4, 2)] %fn1 = block ret func_end @@ -2259,5 +2259,105 @@ 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/disassembler.cc b/src/tint/ir/disassembler.cc index 9211da9bcb..acadde3d94 100644 --- a/src/tint/ir/disassembler.cc +++ b/src/tint/ir/disassembler.cc @@ -20,6 +20,7 @@ #include "src/tint/ir/switch.h" #include "src/tint/ir/terminator.h" #include "src/tint/switch.h" +#include "src/tint/type/type.h" #include "src/tint/utils/scoped_assignment.h" namespace tint::ir { @@ -92,7 +93,9 @@ void Disassembler::Walk(const FlowNode* node) { [&](const ir::Function* f) { TINT_SCOPED_ASSIGNMENT(in_function_, true); - Indent() << "%fn" << GetIdForNode(f) << " = func " << f->name.Name(); + Indent() << "%fn" << GetIdForNode(f) << " = func " << f->name.Name() << "(" + << f->return_type->FriendlyName() << ")"; + if (f->pipeline_stage != Function::PipelineStage::kUndefined) { out_ << " [@" << f->pipeline_stage; @@ -101,6 +104,18 @@ void Disassembler::Walk(const FlowNode* node) { out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] << ")"; } + + if (!f->return_attributes.IsEmpty()) { + out_ << " ra:"; + + for (auto attr : f->return_attributes) { + out_ << " @" << attr; + if (attr == Function::ReturnAttribute::kLocation) { + out_ << "(" << f->return_location.value() << ")"; + } + } + } + out_ << "]"; } out_ << std::endl; diff --git a/src/tint/ir/function.cc b/src/tint/ir/function.cc index 7f73c8964d..a03812e596 100644 --- a/src/tint/ir/function.cc +++ b/src/tint/ir/function.cc @@ -36,4 +36,22 @@ utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStag return out << ""; } +utils::StringStream& operator<<(utils::StringStream& out, Function::ReturnAttribute value) { + switch (value) { + case Function::ReturnAttribute::kLocation: + return out << "location"; + case Function::ReturnAttribute::kFragDepth: + return out << "frag_depth"; + case Function::ReturnAttribute::kSampleMask: + return out << "sample_mask"; + case Function::ReturnAttribute::kPosition: + return out << "position"; + case Function::ReturnAttribute::kInvariant: + return out << "invariant"; + default: + break; + } + return out << ""; +} + } // namespace tint::ir diff --git a/src/tint/ir/function.h b/src/tint/ir/function.h index bbbd8934d4..7393b94b73 100644 --- a/src/tint/ir/function.h +++ b/src/tint/ir/function.h @@ -19,6 +19,7 @@ #include "src/tint/ir/flow_node.h" #include "src/tint/symbol.h" +#include "src/tint/type/type.h" // Forward declarations namespace tint::ir { @@ -43,6 +44,22 @@ class Function : public utils::Castable { kVertex, }; + /// Attributes attached to return types + enum class ReturnAttribute { + /// No return attribute + kNone, + /// Location attribute + kLocation, + /// Builtin Position attribute + kPosition, + /// Builtin FragDepth attribute + kFragDepth, + /// Builtin SampleMask + kSampleMask, + /// Invariant attribute + kInvariant, + }; + /// Constructor Function(); ~Function() override; @@ -56,6 +73,13 @@ class Function : public utils::Castable { /// If this is a `compute` entry point, holds the workgroup size information std::optional> workgroup_size; + /// The function return type + const type::Type* return_type = nullptr; + /// The function return attributes if any + utils::Vector return_attributes; + /// If the return attribute is `kLocation` this stores the location value. + std::optional return_location; + /// The start target is the first block in a function. Block* start_target = nullptr; /// The end target is the end of the function. It is used as the branch target if a return is @@ -64,6 +88,7 @@ class Function : public utils::Castable { }; utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value); +utils::StringStream& operator<<(utils::StringStream& out, Function::ReturnAttribute value); } // namespace tint::ir