[ir] Add function parameters.

This CL adds parameters to functions in the IR. Attributes will be
handled in a later CL.

Bug: tint:1915
Change-Id: I99f9342e95ef957e5117465bb396db0b3822ba7d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/133401
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/tint/BUILD.gn

@@ -1226,6 +1226,8 @@ libtint_source_set("libtint_ir_src") {
     "ir/flow_node.h",
     "ir/function.cc",
     "ir/function.h",
+    "ir/function_param.cc",
+    "ir/function_param.h",
     "ir/function_terminator.cc",
     "ir/function_terminator.h",
     "ir/if.cc",

src/tint/CMakeLists.txt

@@ -742,6 +742,8 @@ if(${TINT_BUILD_IR})
     ir/flow_node.h
     ir/function.cc
     ir/function.h
+    ir/function_param.cc
+    ir/function_param.h
     ir/function_terminator.cc
     ir/function_terminator.h
     ir/if.cc

src/tint/ir/builder.cc

@@ -247,4 +247,8 @@ ir::BlockParam* Builder::BlockParam(const type::Type* type) {
     return ir.values.Create<ir::BlockParam>(type);
 }
 
+ir::FunctionParam* Builder::FunctionParam(const type::Type* type) {
+    return ir.values.Create<ir::FunctionParam>(type);
+}
+
 }  // namespace tint::ir

src/tint/ir/builder.h

@@ -20,12 +20,14 @@
 #include "src/tint/constant/scalar.h"
 #include "src/tint/ir/binary.h"
 #include "src/tint/ir/bitcast.h"
+#include "src/tint/ir/block_param.h"
 #include "src/tint/ir/builtin.h"
 #include "src/tint/ir/constant.h"
 #include "src/tint/ir/construct.h"
 #include "src/tint/ir/convert.h"
 #include "src/tint/ir/discard.h"
 #include "src/tint/ir/function.h"
+#include "src/tint/ir/function_param.h"
 #include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/load.h"
@@ -360,6 +362,11 @@ class Builder {
     /// @returns the value
     ir::BlockParam* BlockParam(const type::Type* type);
 
+    /// Creates a new `FunctionParam`
+    /// @param type the parameter type
+    /// @returns the value
+    ir::FunctionParam* FunctionParam(const type::Type* type);
+
     /// Retrieves the root block for the module, creating if necessary
     /// @returns the root block
     ir::Block* CreateRootBlockIfNeeded();

src/tint/ir/disassembler.cc

@@ -114,8 +114,14 @@ void Disassembler::Walk(const FlowNode* node) {
         [&](const ir::Function* f) {
             TINT_SCOPED_ASSIGNMENT(in_function_, true);
 
-            Indent() << "%fn" << IdOf(f) << " = func " << f->name.Name()
-                     << "():" << f->return_type->FriendlyName();
+            Indent() << "%fn" << IdOf(f) << " = func " << f->name.Name() << "(";
+            for (auto* p : f->params) {
+                if (p != f->params.Front()) {
+                    out_ << ", ";
+                }
+                out_ << "%" << IdOf(p) << ":" << p->type->FriendlyName();
+            }
+            out_ << "):" << f->return_type->FriendlyName();
 
             if (f->pipeline_stage != Function::PipelineStage::kUndefined) {
                 out_ << " [@" << f->pipeline_stage;
@@ -407,6 +413,7 @@ void Disassembler::EmitValue(const Value* val) {
         [&](const ir::BlockParam* p) {
             out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
         },
+        [&](const ir::FunctionParam* p) { out_ << "%" << IdOf(p); },
         [&](Default) { out_ << "Unknown value: " << val->TypeInfo().name; });
 }
 

src/tint/ir/from_program.cc

@@ -335,14 +335,27 @@ class Impl {
         }
         ir_func->return_location = sem->ReturnLocation();
 
+        scopes_.Push();
+        TINT_DEFER(scopes_.Pop());
+
+        utils::Vector<FunctionParam*, 1> params;
+        for (auto* p : ast_func->params) {
+            const auto* param_sem = program_->Sem().Get(p);
+            auto* ty = param_sem->Type()->Clone(clone_ctx_.type_ctx);
+            auto* param = builder_.FunctionParam(ty);
+
+            scopes_.Set(p->name->symbol, param);
+            builder_.ir.SetName(param, p->name->symbol.NameView());
+            params.Push(param);
+        }
+        ir_func->params = std::move(params);
+
         {
             FlowStackScope scope(this, ir_func);
 
             current_flow_block_ = ir_func->start_target;
             EmitBlock(ast_func->body);
 
-            // TODO(dsinclair): Store parameters
-
             // If the branch target has already been set then a `return` was called. Only set in the
             // case where `return` wasn't called.
             BranchToIfNeeded(current_function_->end_target);

src/tint/ir/from_program_binary_test.cc

@@ -830,7 +830,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) {
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:bool):bool {
   %fn2 = block {
   } -> %func_end true # return
 } %func_end

src/tint/ir/from_program_call_test.cc

@@ -77,14 +77,14 @@ TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) {
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():void {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:f32):void {
   %fn2 = block {
   } -> %func_end # return
 } %func_end
 
 %fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
   %fn4 = block {
-    %1:void = call my_func, 6.0f
+    %2:void = call my_func, 6.0f
   } -> %func_end # return
 } %func_end
 

src/tint/ir/from_program_test.cc

@@ -79,6 +79,57 @@ TEST_F(IR_BuilderImplTest, Func) {
 )");
 }
 
+TEST_F(IR_BuilderImplTest, Func_WithParam) {
+    Func("f", utils::Vector{Param("a", ty.u32())}, ty.u32(), utils::Vector{Return("a")});
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    ASSERT_EQ(1u, m->functions.Length());
+
+    auto* f = m->functions[0];
+    ASSERT_NE(f->start_target, nullptr);
+    ASSERT_NE(f->end_target, nullptr);
+
+    EXPECT_EQ(1u, f->start_target->inbound_branches.Length());
+    EXPECT_EQ(1u, f->end_target->inbound_branches.Length());
+
+    EXPECT_EQ(m->functions[0]->pipeline_stage, Function::PipelineStage::kUndefined);
+
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32):u32 {
+  %fn2 = block {
+  } -> %func_end %a # return
+} %func_end
+
+)");
+}
+
+TEST_F(IR_BuilderImplTest, Func_WithMultipleParam) {
+    Func("f", utils::Vector{Param("a", ty.u32()), Param("b", ty.i32()), Param("c", ty.bool_())},
+         ty.void_(), utils::Empty);
+
+    auto m = Build();
+    ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+    ASSERT_EQ(1u, m->functions.Length());
+
+    auto* f = m->functions[0];
+    ASSERT_NE(f->start_target, nullptr);
+    ASSERT_NE(f->end_target, nullptr);
+
+    EXPECT_EQ(1u, f->start_target->inbound_branches.Length());
+    EXPECT_EQ(1u, f->end_target->inbound_branches.Length());
+
+    EXPECT_EQ(m->functions[0]->pipeline_stage, Function::PipelineStage::kUndefined);
+
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32, %b:i32, %c:bool):void {
+  %fn2 = block {
+  } -> %func_end # return
+} %func_end
+
+)");
+}
+
 TEST_F(IR_BuilderImplTest, EntryPoint) {
     Func("f", utils::Empty, ty.void_(), utils::Empty,
          utils::Vector{Stage(ast::PipelineStage::kFragment)});

src/tint/ir/function.h

@@ -19,6 +19,7 @@
 #include <optional>
 
 #include "src/tint/ir/flow_node.h"
+#include "src/tint/ir/function_param.h"
 #include "src/tint/symbol.h"
 #include "src/tint/type/type.h"
 
@@ -88,6 +89,9 @@ class Function : public utils::Castable<Function, FlowNode> {
     /// If the return attribute is `kLocation` this stores the location value.
     std::optional<uint32_t> return_location;
 
+    /// The parameters to the function
+    utils::Vector<FunctionParam*, 1> params;
+
     /// 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

src/tint/ir/function_param.cc

@@ -0,0 +1,25 @@
+// 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/function_param.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::FunctionParam);
+
+namespace tint::ir {
+
+FunctionParam::FunctionParam(const type::Type* ty) : type(ty) {}
+
+FunctionParam::~FunctionParam() = default;
+
+}  // namespace tint::ir

src/tint/ir/function_param.h

@@ -0,0 +1,45 @@
+// 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.
+
+#ifndef SRC_TINT_IR_FUNCTION_PARAM_H_
+#define SRC_TINT_IR_FUNCTION_PARAM_H_
+
+#include "src/tint/ir/value.h"
+#include "src/tint/utils/castable.h"
+
+namespace tint::ir {
+
+/// A function parameter in the IR.
+class FunctionParam : public utils::Castable<FunctionParam, Value> {
+  public:
+    /// Constructor
+    /// @param type the type of the var
+    explicit FunctionParam(const type::Type* type);
+    FunctionParam(const FunctionParam& inst) = delete;
+    FunctionParam(FunctionParam&& inst) = delete;
+    ~FunctionParam() override;
+
+    FunctionParam& operator=(const FunctionParam& inst) = delete;
+    FunctionParam& operator=(FunctionParam&& inst) = delete;
+
+    /// @returns the type of the var
+    const type::Type* Type() const override { return type; }
+
+    /// The type of the parameter
+    const type::Type* type;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_FUNCTION_PARAM_H_