blob: a8faa3cbdcf8d29ed1b05a4e39750b97fbcfcecb [file] [log] [blame]
Austin Engcc2516a2023-10-17 20:57:54 +00001// Copyright 2022 The Dawn & Tint Authors
dan sinclaireee8d882022-10-28 01:22:58 +00002//
Austin Engcc2516a2023-10-17 20:57:54 +00003// Redistribution and use in source and binary forms, with or without
4// modification, are permitted provided that the following conditions are met:
dan sinclaireee8d882022-10-28 01:22:58 +00005//
Austin Engcc2516a2023-10-17 20:57:54 +00006// 1. Redistributions of source code must retain the above copyright notice, this
7// list of conditions and the following disclaimer.
dan sinclaireee8d882022-10-28 01:22:58 +00008//
Austin Engcc2516a2023-10-17 20:57:54 +00009// 2. Redistributions in binary form must reproduce the above copyright notice,
10// this list of conditions and the following disclaimer in the documentation
11// and/or other materials provided with the distribution.
12//
13// 3. Neither the name of the copyright holder nor the names of its
14// contributors may be used to endorse or promote products derived from
15// this software without specific prior written permission.
16//
17// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
18// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
19// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
20// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
21// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
22// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
23// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
24// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
25// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
26// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
dan sinclaireee8d882022-10-28 01:22:58 +000027
dan sinclair97c37272023-07-24 17:11:53 +000028#ifndef SRC_TINT_LANG_CORE_IR_FUNCTION_H_
29#define SRC_TINT_LANG_CORE_IR_FUNCTION_H_
dan sinclaireee8d882022-10-28 01:22:58 +000030
James Priceb298b6a2023-05-04 11:54:19 +000031#include <array>
dan sinclair9d9a3832023-05-03 19:22:31 +000032#include <optional>
dan sinclair24cb8112023-05-18 22:16:08 +000033#include <utility>
dan sinclair9d9a3832023-05-03 19:22:31 +000034
James Price0970b682024-07-04 16:31:11 +000035#include "src/tint/lang/core/io_attributes.h"
dan sinclair97c37272023-07-24 17:11:53 +000036#include "src/tint/lang/core/ir/function_param.h"
dan sinclair97c37272023-07-24 17:11:53 +000037#include "src/tint/lang/core/ir/value.h"
dan sinclair352f8c82023-07-21 00:40:07 +000038#include "src/tint/lang/core/type/type.h"
Ben Claytonacef3102023-11-20 17:09:35 +000039#include "src/tint/utils/containers/const_propagating_ptr.h"
Ben Claytonf848af22023-07-28 16:37:32 +000040#include "src/tint/utils/ice/ice.h"
dan sinclaireee8d882022-10-28 01:22:58 +000041
42// Forward declarations
dan sinclair6f138fe2023-08-15 21:29:34 +000043namespace tint::core::ir {
dan sinclaireee8d882022-10-28 01:22:58 +000044class Block;
dan sinclair09b02ff2023-05-03 22:13:28 +000045class FunctionTerminator;
dan sinclair6f138fe2023-08-15 21:29:34 +000046} // namespace tint::core::ir
dan sinclaireee8d882022-10-28 01:22:58 +000047
dan sinclair6f138fe2023-08-15 21:29:34 +000048namespace tint::core::ir {
dan sinclaireee8d882022-10-28 01:22:58 +000049
50/// An IR representation of a function
dan sinclairbae54e72023-07-28 15:01:54 +000051class Function : public Castable<Function, Value> {
dan sinclaireee8d882022-10-28 01:22:58 +000052 public:
dan sinclair9d9a3832023-05-03 19:22:31 +000053 /// The pipeline stage for an entry point
54 enum class PipelineStage {
55 /// Not a pipeline entry point
56 kUndefined,
57 /// Vertex
58 kCompute,
59 /// Fragment
60 kFragment,
61 /// Vertex
62 kVertex,
63 };
64
dan sinclaireee8d882022-10-28 01:22:58 +000065 /// Constructor
Ben Clayton90c7cc22023-12-01 11:16:00 +000066 Function();
67
68 /// Constructor
dan sinclairc9923d22023-05-15 23:44:04 +000069 /// @param rt the function return type
dan sinclairf59547f2023-05-16 14:43:45 +000070 /// @param stage the function stage
71 /// @param wg_size the workgroup_size
dan sinclaircedcdf32023-08-10 02:39:48 +000072 Function(const core::type::Type* rt,
dan sinclairf59547f2023-05-16 14:43:45 +000073 PipelineStage stage = PipelineStage::kUndefined,
74 std::optional<std::array<uint32_t, 3>> wg_size = {});
dan sinclaireee8d882022-10-28 01:22:58 +000075 ~Function() override;
76
dan sinclair16cb4bd2023-09-21 08:52:11 +000077 /// @copydoc Instruction::Clone()
78 Function* Clone(CloneContext& ctx) override;
79
dan sinclair24cb8112023-05-18 22:16:08 +000080 /// Sets the function stage
81 /// @param stage the stage to set
82 void SetStage(PipelineStage stage) { pipeline_stage_ = stage; }
dan sinclair9d9a3832023-05-03 19:22:31 +000083
dan sinclair24cb8112023-05-18 22:16:08 +000084 /// @returns the function pipeline stage
Ben Claytonead8a042023-11-18 09:45:32 +000085 PipelineStage Stage() const { return pipeline_stage_; }
dan sinclair69bb5dd2023-05-03 21:31:51 +000086
dan sinclair24cb8112023-05-18 22:16:08 +000087 /// Sets the workgroup size
88 /// @param x the x size
89 /// @param y the y size
90 /// @param z the z size
91 void SetWorkgroupSize(uint32_t x, uint32_t y, uint32_t z) { workgroup_size_ = {x, y, z}; }
dan sinclair84d750e2023-05-18 08:50:34 +000092
Ben Claytond8bde2c2024-05-28 13:46:59 +000093 /// Sets the workgroup size
94 /// @param size the new size
95 void SetWorkgroupSize(std::array<uint32_t, 3> size) { workgroup_size_ = size; }
96
James Price9e819bf2023-06-28 12:13:41 +000097 /// Clears the workgroup size.
98 void ClearWorkgroupSize() { workgroup_size_ = {}; }
99
dan sinclair24cb8112023-05-18 22:16:08 +0000100 /// @returns the workgroup size information
Ben Claytonead8a042023-11-18 09:45:32 +0000101 std::optional<std::array<uint32_t, 3>> WorkgroupSize() const { return workgroup_size_; }
dan sinclair24cb8112023-05-18 22:16:08 +0000102
Ben Clayton90c7cc22023-12-01 11:16:00 +0000103 /// @param type the return type for the function
104 void SetReturnType(const core::type::Type* type) { return_.type = type; }
105
dan sinclair24cb8112023-05-18 22:16:08 +0000106 /// @returns the return type for the function
Ben Claytonead8a042023-11-18 09:45:32 +0000107 const core::type::Type* ReturnType() const { return return_.type; }
dan sinclair24cb8112023-05-18 22:16:08 +0000108
James Price0970b682024-07-04 16:31:11 +0000109 /// Sets the return IO attributes.
110 /// @param attrs the attributes
111 void SetReturnAttributes(const IOAttributes& attrs) { return_.attributes = attrs; }
112 /// @returns the return IO attributes
113 const IOAttributes& ReturnAttributes() const { return return_.attributes; }
114
dan sinclair24cb8112023-05-18 22:16:08 +0000115 /// Sets the return attributes
dan sinclair212959b2023-05-30 19:55:08 +0000116 /// @param builtin the builtin to set
Ben Clayton3b0e5df2024-01-03 16:49:21 +0000117 void SetReturnBuiltin(BuiltinValue builtin) {
James Price0970b682024-07-04 16:31:11 +0000118 TINT_ASSERT(!return_.attributes.builtin.has_value());
119 return_.attributes.builtin = builtin;
dan sinclair24cb8112023-05-18 22:16:08 +0000120 }
dan sinclair212959b2023-05-30 19:55:08 +0000121 /// @returns the return builtin attribute
James Price0970b682024-07-04 16:31:11 +0000122 std::optional<BuiltinValue> ReturnBuiltin() const { return return_.attributes.builtin; }
Ben Claytone8b7b022024-01-03 16:49:21 +0000123
James Price963cfb02024-07-04 21:51:16 +0000124 /// Sets the return location.
125 /// @param loc the optional location to set
126 void SetReturnLocation(std::optional<uint32_t> loc) { return_.attributes.location = loc; }
127
James Price0ca73db2024-07-04 21:51:23 +0000128 /// @returns the return location
129 std::optional<uint32_t> ReturnLocation() const { return return_.attributes.location; }
130
James Price963cfb02024-07-04 21:51:16 +0000131 /// Sets the return interpolation.
132 /// @param interp the optional interpolation
133 void SetReturnInterpolation(std::optional<core::Interpolation> interp) {
James Price0970b682024-07-04 16:31:11 +0000134 return_.attributes.interpolation = interp;
dan sinclair212959b2023-05-30 19:55:08 +0000135 }
Ben Claytone8b7b022024-01-03 16:49:21 +0000136
James Price0ca73db2024-07-04 21:51:23 +0000137 /// @returns the return interpolation
138 std::optional<Interpolation> ReturnInterpolation() const {
139 return return_.attributes.interpolation;
James Price0970b682024-07-04 16:31:11 +0000140 }
Ben Claytone8b7b022024-01-03 16:49:21 +0000141
dan sinclair212959b2023-05-30 19:55:08 +0000142 /// Sets the return as invariant
143 /// @param val the invariant value to set
James Price0970b682024-07-04 16:31:11 +0000144 void SetReturnInvariant(bool val) { return_.attributes.invariant = val; }
Ben Claytone8b7b022024-01-03 16:49:21 +0000145
dan sinclair212959b2023-05-30 19:55:08 +0000146 /// @returns the return invariant value
James Price0970b682024-07-04 16:31:11 +0000147 bool ReturnInvariant() const { return return_.attributes.invariant; }
dan sinclair24cb8112023-05-18 22:16:08 +0000148
149 /// Sets the function parameters
Ben Claytonb768af42023-06-09 23:31:16 +0000150 /// @param params the function parameters
dan sinclairbae54e72023-07-28 15:01:54 +0000151 void SetParams(VectorRef<FunctionParam*> params);
Ben Claytonb768af42023-06-09 23:31:16 +0000152
153 /// Sets the function parameters
154 /// @param params the function parameters
155 void SetParams(std::initializer_list<FunctionParam*> params);
156
James Pricede601a82024-05-07 01:34:10 +0000157 /// Appends a new function parameter.
158 /// @param param the function parameter to append
159 void AppendParam(FunctionParam* param);
160
dan sinclair24cb8112023-05-18 22:16:08 +0000161 /// @returns the function parameters
dan sinclairbae54e72023-07-28 15:01:54 +0000162 const VectorRef<FunctionParam*> Params() { return params_; }
dan sinclair24cb8112023-05-18 22:16:08 +0000163
Ben Claytonead8a042023-11-18 09:45:32 +0000164 /// @returns the function parameters
165 VectorRef<const FunctionParam*> Params() const { return params_; }
166
Ben Clayton5acd44c2023-06-22 13:38:30 +0000167 /// Sets the root block for the function
168 /// @param target the root block
169 void SetBlock(Block* target) {
Ben Claytonf848af22023-07-28 16:37:32 +0000170 TINT_ASSERT(target != nullptr);
Ben Clayton5acd44c2023-06-22 13:38:30 +0000171 block_ = target;
dan sinclaireefa7fe2023-06-05 15:16:53 +0000172 }
Ben Claytonead8a042023-11-18 09:45:32 +0000173
Ben Clayton5acd44c2023-06-22 13:38:30 +0000174 /// @returns the function root block
175 ir::Block* Block() { return block_; }
dan sinclair24cb8112023-05-18 22:16:08 +0000176
Ben Claytonead8a042023-11-18 09:45:32 +0000177 /// @returns the function root block
178 const ir::Block* Block() const { return block_; }
179
James Pricec00c5692023-10-03 14:39:27 +0000180 /// Destroys the function and all of its instructions.
181 void Destroy() override;
182
dan sinclair24cb8112023-05-18 22:16:08 +0000183 private:
Ben Clayton90c7cc22023-12-01 11:16:00 +0000184 PipelineStage pipeline_stage_ = PipelineStage::kUndefined;
dan sinclair24cb8112023-05-18 22:16:08 +0000185 std::optional<std::array<uint32_t, 3>> workgroup_size_;
186
dan sinclair212959b2023-05-30 19:55:08 +0000187 struct {
dan sinclaircedcdf32023-08-10 02:39:48 +0000188 const core::type::Type* type = nullptr;
James Price0970b682024-07-04 16:31:11 +0000189 IOAttributes attributes = {};
dan sinclair212959b2023-05-30 19:55:08 +0000190 } return_;
dan sinclair24cb8112023-05-18 22:16:08 +0000191
dan sinclairbae54e72023-07-28 15:01:54 +0000192 Vector<FunctionParam*, 1> params_;
Ben Claytonacef3102023-11-20 17:09:35 +0000193 ConstPropagatingPtr<ir::Block> block_;
dan sinclaireee8d882022-10-28 01:22:58 +0000194};
195
Ben Clayton68919602023-07-28 22:51:18 +0000196/// @param value the enum value
197/// @returns the string for the given enum value
198std::string_view ToString(Function::PipelineStage value);
199
200/// @param out the stream to write to
201/// @param value the Function::PipelineStage
202/// @returns @p out so calls can be chained
203template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>>
204auto& operator<<(STREAM& out, Function::PipelineStage value) {
205 return out << ToString(value);
206}
207
dan sinclair6f138fe2023-08-15 21:29:34 +0000208} // namespace tint::core::ir
dan sinclaireee8d882022-10-28 01:22:58 +0000209
dan sinclair97c37272023-07-24 17:11:53 +0000210#endif // SRC_TINT_LANG_CORE_IR_FUNCTION_H_