Austin Eng | cc2516a | 2023-10-17 20:57:54 +0000 | [diff] [blame] | 1 | // Copyright 2022 The Dawn & Tint Authors |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 2 | // |
Austin Eng | cc2516a | 2023-10-17 20:57:54 +0000 | [diff] [blame] | 3 | // Redistribution and use in source and binary forms, with or without |
| 4 | // modification, are permitted provided that the following conditions are met: |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 5 | // |
Austin Eng | cc2516a | 2023-10-17 20:57:54 +0000 | [diff] [blame] | 6 | // 1. Redistributions of source code must retain the above copyright notice, this |
| 7 | // list of conditions and the following disclaimer. |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 8 | // |
Austin Eng | cc2516a | 2023-10-17 20:57:54 +0000 | [diff] [blame] | 9 | // 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 sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 27 | |
dan sinclair | 97c3727 | 2023-07-24 17:11:53 +0000 | [diff] [blame] | 28 | #ifndef SRC_TINT_LANG_CORE_IR_FUNCTION_H_ |
| 29 | #define SRC_TINT_LANG_CORE_IR_FUNCTION_H_ |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 30 | |
James Price | b298b6a | 2023-05-04 11:54:19 +0000 | [diff] [blame] | 31 | #include <array> |
dan sinclair | 9d9a383 | 2023-05-03 19:22:31 +0000 | [diff] [blame] | 32 | #include <optional> |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 33 | #include <utility> |
dan sinclair | 9d9a383 | 2023-05-03 19:22:31 +0000 | [diff] [blame] | 34 | |
James Price | 0970b68 | 2024-07-04 16:31:11 +0000 | [diff] [blame] | 35 | #include "src/tint/lang/core/io_attributes.h" |
dan sinclair | 97c3727 | 2023-07-24 17:11:53 +0000 | [diff] [blame] | 36 | #include "src/tint/lang/core/ir/function_param.h" |
dan sinclair | 97c3727 | 2023-07-24 17:11:53 +0000 | [diff] [blame] | 37 | #include "src/tint/lang/core/ir/value.h" |
dan sinclair | 352f8c8 | 2023-07-21 00:40:07 +0000 | [diff] [blame] | 38 | #include "src/tint/lang/core/type/type.h" |
Ben Clayton | acef310 | 2023-11-20 17:09:35 +0000 | [diff] [blame] | 39 | #include "src/tint/utils/containers/const_propagating_ptr.h" |
Ben Clayton | f848af2 | 2023-07-28 16:37:32 +0000 | [diff] [blame] | 40 | #include "src/tint/utils/ice/ice.h" |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 41 | |
| 42 | // Forward declarations |
dan sinclair | 6f138fe | 2023-08-15 21:29:34 +0000 | [diff] [blame] | 43 | namespace tint::core::ir { |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 44 | class Block; |
dan sinclair | 09b02ff | 2023-05-03 22:13:28 +0000 | [diff] [blame] | 45 | class FunctionTerminator; |
dan sinclair | 6f138fe | 2023-08-15 21:29:34 +0000 | [diff] [blame] | 46 | } // namespace tint::core::ir |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 47 | |
dan sinclair | 6f138fe | 2023-08-15 21:29:34 +0000 | [diff] [blame] | 48 | namespace tint::core::ir { |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 49 | |
| 50 | /// An IR representation of a function |
dan sinclair | bae54e7 | 2023-07-28 15:01:54 +0000 | [diff] [blame] | 51 | class Function : public Castable<Function, Value> { |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 52 | public: |
dan sinclair | 9d9a383 | 2023-05-03 19:22:31 +0000 | [diff] [blame] | 53 | /// 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 sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 65 | /// Constructor |
Ben Clayton | 90c7cc2 | 2023-12-01 11:16:00 +0000 | [diff] [blame] | 66 | Function(); |
| 67 | |
| 68 | /// Constructor |
dan sinclair | c9923d2 | 2023-05-15 23:44:04 +0000 | [diff] [blame] | 69 | /// @param rt the function return type |
dan sinclair | f59547f | 2023-05-16 14:43:45 +0000 | [diff] [blame] | 70 | /// @param stage the function stage |
| 71 | /// @param wg_size the workgroup_size |
dan sinclair | cedcdf3 | 2023-08-10 02:39:48 +0000 | [diff] [blame] | 72 | Function(const core::type::Type* rt, |
dan sinclair | f59547f | 2023-05-16 14:43:45 +0000 | [diff] [blame] | 73 | PipelineStage stage = PipelineStage::kUndefined, |
| 74 | std::optional<std::array<uint32_t, 3>> wg_size = {}); |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 75 | ~Function() override; |
| 76 | |
dan sinclair | 16cb4bd | 2023-09-21 08:52:11 +0000 | [diff] [blame] | 77 | /// @copydoc Instruction::Clone() |
| 78 | Function* Clone(CloneContext& ctx) override; |
| 79 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 80 | /// Sets the function stage |
| 81 | /// @param stage the stage to set |
| 82 | void SetStage(PipelineStage stage) { pipeline_stage_ = stage; } |
dan sinclair | 9d9a383 | 2023-05-03 19:22:31 +0000 | [diff] [blame] | 83 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 84 | /// @returns the function pipeline stage |
Ben Clayton | ead8a04 | 2023-11-18 09:45:32 +0000 | [diff] [blame] | 85 | PipelineStage Stage() const { return pipeline_stage_; } |
dan sinclair | 69bb5dd | 2023-05-03 21:31:51 +0000 | [diff] [blame] | 86 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 87 | /// 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 sinclair | 84d750e | 2023-05-18 08:50:34 +0000 | [diff] [blame] | 92 | |
Ben Clayton | d8bde2c | 2024-05-28 13:46:59 +0000 | [diff] [blame] | 93 | /// Sets the workgroup size |
| 94 | /// @param size the new size |
| 95 | void SetWorkgroupSize(std::array<uint32_t, 3> size) { workgroup_size_ = size; } |
| 96 | |
James Price | 9e819bf | 2023-06-28 12:13:41 +0000 | [diff] [blame] | 97 | /// Clears the workgroup size. |
| 98 | void ClearWorkgroupSize() { workgroup_size_ = {}; } |
| 99 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 100 | /// @returns the workgroup size information |
Ben Clayton | ead8a04 | 2023-11-18 09:45:32 +0000 | [diff] [blame] | 101 | std::optional<std::array<uint32_t, 3>> WorkgroupSize() const { return workgroup_size_; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 102 | |
Ben Clayton | 90c7cc2 | 2023-12-01 11:16:00 +0000 | [diff] [blame] | 103 | /// @param type the return type for the function |
| 104 | void SetReturnType(const core::type::Type* type) { return_.type = type; } |
| 105 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 106 | /// @returns the return type for the function |
Ben Clayton | ead8a04 | 2023-11-18 09:45:32 +0000 | [diff] [blame] | 107 | const core::type::Type* ReturnType() const { return return_.type; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 108 | |
James Price | 0970b68 | 2024-07-04 16:31:11 +0000 | [diff] [blame] | 109 | /// 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 sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 115 | /// Sets the return attributes |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 116 | /// @param builtin the builtin to set |
Ben Clayton | 3b0e5df | 2024-01-03 16:49:21 +0000 | [diff] [blame] | 117 | void SetReturnBuiltin(BuiltinValue builtin) { |
James Price | 0970b68 | 2024-07-04 16:31:11 +0000 | [diff] [blame] | 118 | TINT_ASSERT(!return_.attributes.builtin.has_value()); |
| 119 | return_.attributes.builtin = builtin; |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 120 | } |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 121 | /// @returns the return builtin attribute |
James Price | 0970b68 | 2024-07-04 16:31:11 +0000 | [diff] [blame] | 122 | std::optional<BuiltinValue> ReturnBuiltin() const { return return_.attributes.builtin; } |
Ben Clayton | e8b7b02 | 2024-01-03 16:49:21 +0000 | [diff] [blame] | 123 | |
James Price | 963cfb0 | 2024-07-04 21:51:16 +0000 | [diff] [blame] | 124 | /// 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 Price | 0ca73db | 2024-07-04 21:51:23 +0000 | [diff] [blame] | 128 | /// @returns the return location |
| 129 | std::optional<uint32_t> ReturnLocation() const { return return_.attributes.location; } |
| 130 | |
James Price | 963cfb0 | 2024-07-04 21:51:16 +0000 | [diff] [blame] | 131 | /// Sets the return interpolation. |
| 132 | /// @param interp the optional interpolation |
| 133 | void SetReturnInterpolation(std::optional<core::Interpolation> interp) { |
James Price | 0970b68 | 2024-07-04 16:31:11 +0000 | [diff] [blame] | 134 | return_.attributes.interpolation = interp; |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 135 | } |
Ben Clayton | e8b7b02 | 2024-01-03 16:49:21 +0000 | [diff] [blame] | 136 | |
James Price | 0ca73db | 2024-07-04 21:51:23 +0000 | [diff] [blame] | 137 | /// @returns the return interpolation |
| 138 | std::optional<Interpolation> ReturnInterpolation() const { |
| 139 | return return_.attributes.interpolation; |
James Price | 0970b68 | 2024-07-04 16:31:11 +0000 | [diff] [blame] | 140 | } |
Ben Clayton | e8b7b02 | 2024-01-03 16:49:21 +0000 | [diff] [blame] | 141 | |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 142 | /// Sets the return as invariant |
| 143 | /// @param val the invariant value to set |
James Price | 0970b68 | 2024-07-04 16:31:11 +0000 | [diff] [blame] | 144 | void SetReturnInvariant(bool val) { return_.attributes.invariant = val; } |
Ben Clayton | e8b7b02 | 2024-01-03 16:49:21 +0000 | [diff] [blame] | 145 | |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 146 | /// @returns the return invariant value |
James Price | 0970b68 | 2024-07-04 16:31:11 +0000 | [diff] [blame] | 147 | bool ReturnInvariant() const { return return_.attributes.invariant; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 148 | |
| 149 | /// Sets the function parameters |
Ben Clayton | b768af4 | 2023-06-09 23:31:16 +0000 | [diff] [blame] | 150 | /// @param params the function parameters |
dan sinclair | bae54e7 | 2023-07-28 15:01:54 +0000 | [diff] [blame] | 151 | void SetParams(VectorRef<FunctionParam*> params); |
Ben Clayton | b768af4 | 2023-06-09 23:31:16 +0000 | [diff] [blame] | 152 | |
| 153 | /// Sets the function parameters |
| 154 | /// @param params the function parameters |
| 155 | void SetParams(std::initializer_list<FunctionParam*> params); |
| 156 | |
James Price | de601a8 | 2024-05-07 01:34:10 +0000 | [diff] [blame] | 157 | /// Appends a new function parameter. |
| 158 | /// @param param the function parameter to append |
| 159 | void AppendParam(FunctionParam* param); |
| 160 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 161 | /// @returns the function parameters |
dan sinclair | bae54e7 | 2023-07-28 15:01:54 +0000 | [diff] [blame] | 162 | const VectorRef<FunctionParam*> Params() { return params_; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 163 | |
Ben Clayton | ead8a04 | 2023-11-18 09:45:32 +0000 | [diff] [blame] | 164 | /// @returns the function parameters |
| 165 | VectorRef<const FunctionParam*> Params() const { return params_; } |
| 166 | |
Ben Clayton | 5acd44c | 2023-06-22 13:38:30 +0000 | [diff] [blame] | 167 | /// Sets the root block for the function |
| 168 | /// @param target the root block |
| 169 | void SetBlock(Block* target) { |
Ben Clayton | f848af2 | 2023-07-28 16:37:32 +0000 | [diff] [blame] | 170 | TINT_ASSERT(target != nullptr); |
Ben Clayton | 5acd44c | 2023-06-22 13:38:30 +0000 | [diff] [blame] | 171 | block_ = target; |
dan sinclair | eefa7fe | 2023-06-05 15:16:53 +0000 | [diff] [blame] | 172 | } |
Ben Clayton | ead8a04 | 2023-11-18 09:45:32 +0000 | [diff] [blame] | 173 | |
Ben Clayton | 5acd44c | 2023-06-22 13:38:30 +0000 | [diff] [blame] | 174 | /// @returns the function root block |
| 175 | ir::Block* Block() { return block_; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 176 | |
Ben Clayton | ead8a04 | 2023-11-18 09:45:32 +0000 | [diff] [blame] | 177 | /// @returns the function root block |
| 178 | const ir::Block* Block() const { return block_; } |
| 179 | |
James Price | c00c569 | 2023-10-03 14:39:27 +0000 | [diff] [blame] | 180 | /// Destroys the function and all of its instructions. |
| 181 | void Destroy() override; |
| 182 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 183 | private: |
Ben Clayton | 90c7cc2 | 2023-12-01 11:16:00 +0000 | [diff] [blame] | 184 | PipelineStage pipeline_stage_ = PipelineStage::kUndefined; |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 185 | std::optional<std::array<uint32_t, 3>> workgroup_size_; |
| 186 | |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 187 | struct { |
dan sinclair | cedcdf3 | 2023-08-10 02:39:48 +0000 | [diff] [blame] | 188 | const core::type::Type* type = nullptr; |
James Price | 0970b68 | 2024-07-04 16:31:11 +0000 | [diff] [blame] | 189 | IOAttributes attributes = {}; |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 190 | } return_; |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 191 | |
dan sinclair | bae54e7 | 2023-07-28 15:01:54 +0000 | [diff] [blame] | 192 | Vector<FunctionParam*, 1> params_; |
Ben Clayton | acef310 | 2023-11-20 17:09:35 +0000 | [diff] [blame] | 193 | ConstPropagatingPtr<ir::Block> block_; |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 194 | }; |
| 195 | |
Ben Clayton | 6891960 | 2023-07-28 22:51:18 +0000 | [diff] [blame] | 196 | /// @param value the enum value |
| 197 | /// @returns the string for the given enum value |
| 198 | std::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 |
| 203 | template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>> |
| 204 | auto& operator<<(STREAM& out, Function::PipelineStage value) { |
| 205 | return out << ToString(value); |
| 206 | } |
| 207 | |
dan sinclair | 6f138fe | 2023-08-15 21:29:34 +0000 | [diff] [blame] | 208 | } // namespace tint::core::ir |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 209 | |
dan sinclair | 97c3727 | 2023-07-24 17:11:53 +0000 | [diff] [blame] | 210 | #endif // SRC_TINT_LANG_CORE_IR_FUNCTION_H_ |