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 | |
dan sinclair | 97c3727 | 2023-07-24 17:11:53 +0000 | [diff] [blame] | 35 | #include "src/tint/lang/core/ir/function_param.h" |
| 36 | #include "src/tint/lang/core/ir/location.h" |
| 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 | f848af2 | 2023-07-28 16:37:32 +0000 | [diff] [blame] | 39 | #include "src/tint/utils/ice/ice.h" |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 40 | |
| 41 | // Forward declarations |
dan sinclair | 6f138fe | 2023-08-15 21:29:34 +0000 | [diff] [blame] | 42 | namespace tint::core::ir { |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 43 | class Block; |
dan sinclair | 09b02ff | 2023-05-03 22:13:28 +0000 | [diff] [blame] | 44 | class FunctionTerminator; |
dan sinclair | 6f138fe | 2023-08-15 21:29:34 +0000 | [diff] [blame] | 45 | } // namespace tint::core::ir |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 46 | |
dan sinclair | 6f138fe | 2023-08-15 21:29:34 +0000 | [diff] [blame] | 47 | namespace tint::core::ir { |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 48 | |
| 49 | /// An IR representation of a function |
dan sinclair | bae54e7 | 2023-07-28 15:01:54 +0000 | [diff] [blame] | 50 | class Function : public Castable<Function, Value> { |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 51 | public: |
dan sinclair | 9d9a383 | 2023-05-03 19:22:31 +0000 | [diff] [blame] | 52 | /// The pipeline stage for an entry point |
| 53 | enum class PipelineStage { |
| 54 | /// Not a pipeline entry point |
| 55 | kUndefined, |
| 56 | /// Vertex |
| 57 | kCompute, |
| 58 | /// Fragment |
| 59 | kFragment, |
| 60 | /// Vertex |
| 61 | kVertex, |
| 62 | }; |
| 63 | |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 64 | /// Builtin attached to return types |
| 65 | enum class ReturnBuiltin { |
dan sinclair | 69bb5dd | 2023-05-03 21:31:51 +0000 | [diff] [blame] | 66 | /// Builtin Position attribute |
| 67 | kPosition, |
| 68 | /// Builtin FragDepth attribute |
| 69 | kFragDepth, |
| 70 | /// Builtin SampleMask |
| 71 | kSampleMask, |
dan sinclair | 69bb5dd | 2023-05-03 21:31:51 +0000 | [diff] [blame] | 72 | }; |
| 73 | |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 74 | /// Constructor |
dan sinclair | c9923d2 | 2023-05-15 23:44:04 +0000 | [diff] [blame] | 75 | /// @param rt the function return type |
dan sinclair | f59547f | 2023-05-16 14:43:45 +0000 | [diff] [blame] | 76 | /// @param stage the function stage |
| 77 | /// @param wg_size the workgroup_size |
dan sinclair | cedcdf3 | 2023-08-10 02:39:48 +0000 | [diff] [blame] | 78 | Function(const core::type::Type* rt, |
dan sinclair | f59547f | 2023-05-16 14:43:45 +0000 | [diff] [blame] | 79 | PipelineStage stage = PipelineStage::kUndefined, |
| 80 | std::optional<std::array<uint32_t, 3>> wg_size = {}); |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 81 | ~Function() override; |
| 82 | |
dan sinclair | 16cb4bd | 2023-09-21 08:52:11 +0000 | [diff] [blame] | 83 | /// @copydoc Instruction::Clone() |
| 84 | Function* Clone(CloneContext& ctx) override; |
| 85 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 86 | /// Sets the function stage |
| 87 | /// @param stage the stage to set |
| 88 | void SetStage(PipelineStage stage) { pipeline_stage_ = stage; } |
dan sinclair | 9d9a383 | 2023-05-03 19:22:31 +0000 | [diff] [blame] | 89 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 90 | /// @returns the function pipeline stage |
dan sinclair | a5268df | 2023-06-08 21:39:53 +0000 | [diff] [blame] | 91 | PipelineStage Stage() { return pipeline_stage_; } |
dan sinclair | 69bb5dd | 2023-05-03 21:31:51 +0000 | [diff] [blame] | 92 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 93 | /// Sets the workgroup size |
| 94 | /// @param x the x size |
| 95 | /// @param y the y size |
| 96 | /// @param z the z size |
| 97 | 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] | 98 | |
James Price | 9e819bf | 2023-06-28 12:13:41 +0000 | [diff] [blame] | 99 | /// Clears the workgroup size. |
| 100 | void ClearWorkgroupSize() { workgroup_size_ = {}; } |
| 101 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 102 | /// @returns the workgroup size information |
dan sinclair | a5268df | 2023-06-08 21:39:53 +0000 | [diff] [blame] | 103 | std::optional<std::array<uint32_t, 3>> WorkgroupSize() { return workgroup_size_; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 104 | |
| 105 | /// @returns the return type for the function |
dan sinclair | cedcdf3 | 2023-08-10 02:39:48 +0000 | [diff] [blame] | 106 | const core::type::Type* ReturnType() { return return_.type; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 107 | |
| 108 | /// Sets the return attributes |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 109 | /// @param builtin the builtin to set |
| 110 | void SetReturnBuiltin(ReturnBuiltin builtin) { |
Ben Clayton | f848af2 | 2023-07-28 16:37:32 +0000 | [diff] [blame] | 111 | TINT_ASSERT(!return_.builtin.has_value()); |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 112 | return_.builtin = builtin; |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 113 | } |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 114 | /// @returns the return builtin attribute |
dan sinclair | a5268df | 2023-06-08 21:39:53 +0000 | [diff] [blame] | 115 | std::optional<enum ReturnBuiltin> ReturnBuiltin() { return return_.builtin; } |
James Price | 9e819bf | 2023-06-28 12:13:41 +0000 | [diff] [blame] | 116 | /// Clears the return builtin attribute. |
| 117 | void ClearReturnBuiltin() { return_.builtin = {}; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 118 | |
| 119 | /// Sets the return location |
| 120 | /// @param loc the location to set |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 121 | /// @param interp the interpolation |
Ben Clayton | cd52f38 | 2023-08-07 13:11:08 +0000 | [diff] [blame] | 122 | void SetReturnLocation(uint32_t loc, std::optional<core::Interpolation> interp) { |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 123 | return_.location = {loc, interp}; |
| 124 | } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 125 | /// @returns the return location |
dan sinclair | a5268df | 2023-06-08 21:39:53 +0000 | [diff] [blame] | 126 | std::optional<Location> ReturnLocation() { return return_.location; } |
James Price | 9e819bf | 2023-06-28 12:13:41 +0000 | [diff] [blame] | 127 | /// Clears the return location attribute. |
| 128 | void ClearReturnLocation() { return_.location = {}; } |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 129 | |
| 130 | /// Sets the return as invariant |
| 131 | /// @param val the invariant value to set |
| 132 | void SetReturnInvariant(bool val) { return_.invariant = val; } |
| 133 | /// @returns the return invariant value |
dan sinclair | a5268df | 2023-06-08 21:39:53 +0000 | [diff] [blame] | 134 | bool ReturnInvariant() { return return_.invariant; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 135 | |
| 136 | /// Sets the function parameters |
Ben Clayton | b768af4 | 2023-06-09 23:31:16 +0000 | [diff] [blame] | 137 | /// @param params the function parameters |
dan sinclair | bae54e7 | 2023-07-28 15:01:54 +0000 | [diff] [blame] | 138 | void SetParams(VectorRef<FunctionParam*> params); |
Ben Clayton | b768af4 | 2023-06-09 23:31:16 +0000 | [diff] [blame] | 139 | |
| 140 | /// Sets the function parameters |
| 141 | /// @param params the function parameters |
| 142 | void SetParams(std::initializer_list<FunctionParam*> params); |
| 143 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 144 | /// @returns the function parameters |
dan sinclair | bae54e7 | 2023-07-28 15:01:54 +0000 | [diff] [blame] | 145 | const VectorRef<FunctionParam*> Params() { return params_; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 146 | |
Ben Clayton | 5acd44c | 2023-06-22 13:38:30 +0000 | [diff] [blame] | 147 | /// Sets the root block for the function |
| 148 | /// @param target the root block |
| 149 | void SetBlock(Block* target) { |
Ben Clayton | f848af2 | 2023-07-28 16:37:32 +0000 | [diff] [blame] | 150 | TINT_ASSERT(target != nullptr); |
Ben Clayton | 5acd44c | 2023-06-22 13:38:30 +0000 | [diff] [blame] | 151 | block_ = target; |
dan sinclair | eefa7fe | 2023-06-05 15:16:53 +0000 | [diff] [blame] | 152 | } |
Ben Clayton | 5acd44c | 2023-06-22 13:38:30 +0000 | [diff] [blame] | 153 | /// @returns the function root block |
| 154 | ir::Block* Block() { return block_; } |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 155 | |
James Price | c00c569 | 2023-10-03 14:39:27 +0000 | [diff] [blame] | 156 | /// Destroys the function and all of its instructions. |
| 157 | void Destroy() override; |
| 158 | |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 159 | private: |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 160 | PipelineStage pipeline_stage_; |
| 161 | std::optional<std::array<uint32_t, 3>> workgroup_size_; |
| 162 | |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 163 | struct { |
dan sinclair | cedcdf3 | 2023-08-10 02:39:48 +0000 | [diff] [blame] | 164 | const core::type::Type* type = nullptr; |
dan sinclair | 212959b | 2023-05-30 19:55:08 +0000 | [diff] [blame] | 165 | std::optional<enum ReturnBuiltin> builtin; |
| 166 | std::optional<Location> location; |
| 167 | bool invariant = false; |
| 168 | } return_; |
dan sinclair | 24cb811 | 2023-05-18 22:16:08 +0000 | [diff] [blame] | 169 | |
dan sinclair | bae54e7 | 2023-07-28 15:01:54 +0000 | [diff] [blame] | 170 | Vector<FunctionParam*, 1> params_; |
Ben Clayton | 5acd44c | 2023-06-22 13:38:30 +0000 | [diff] [blame] | 171 | ir::Block* block_ = nullptr; |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 172 | }; |
| 173 | |
Ben Clayton | 6891960 | 2023-07-28 22:51:18 +0000 | [diff] [blame] | 174 | /// @param value the enum value |
| 175 | /// @returns the string for the given enum value |
| 176 | std::string_view ToString(Function::PipelineStage value); |
| 177 | |
| 178 | /// @param out the stream to write to |
| 179 | /// @param value the Function::PipelineStage |
| 180 | /// @returns @p out so calls can be chained |
| 181 | template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>> |
| 182 | auto& operator<<(STREAM& out, Function::PipelineStage value) { |
| 183 | return out << ToString(value); |
| 184 | } |
| 185 | |
| 186 | /// @param value the enum value |
| 187 | /// @returns the string for the given enum value |
| 188 | std::string_view ToString(enum Function::ReturnBuiltin value); |
| 189 | |
| 190 | /// @param out the stream to write to |
| 191 | /// @param value the Function::ReturnBuiltin |
| 192 | /// @returns @p out so calls can be chained |
| 193 | template <typename STREAM, typename = traits::EnableIfIsOStream<STREAM>> |
| 194 | auto& operator<<(STREAM& out, enum Function::ReturnBuiltin value) { |
| 195 | return out << ToString(value); |
| 196 | } |
dan sinclair | 9d9a383 | 2023-05-03 19:22:31 +0000 | [diff] [blame] | 197 | |
dan sinclair | 6f138fe | 2023-08-15 21:29:34 +0000 | [diff] [blame] | 198 | } // namespace tint::core::ir |
dan sinclair | eee8d88 | 2022-10-28 01:22:58 +0000 | [diff] [blame] | 199 | |
dan sinclair | 97c3727 | 2023-07-24 17:11:53 +0000 | [diff] [blame] | 200 | #endif // SRC_TINT_LANG_CORE_IR_FUNCTION_H_ |