|
| 1 | +// Copyright 2024 The Dawn & Tint Authors |
| 2 | +// |
| 3 | +// Redistribution and use in source and binary forms, with or without |
| 4 | +// modification, are permitted provided that the following conditions are met: |
| 5 | +// |
| 6 | +// 1. Redistributions of source code must retain the above copyright notice, this |
| 7 | +// list of conditions and the following disclaimer. |
| 8 | +// |
| 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. |
| 27 | + |
| 28 | +#ifndef SRC_TINT_LANG_CORE_IR_REFERENCED_MODULE_DECLS_H_ |
| 29 | +#define SRC_TINT_LANG_CORE_IR_REFERENCED_MODULE_DECLS_H_ |
| 30 | + |
| 31 | +#include "src/tint/lang/core/ir/control_instruction.h" |
| 32 | +#include "src/tint/lang/core/ir/instruction.h" |
| 33 | +#include "src/tint/lang/core/ir/let.h" |
| 34 | +#include "src/tint/lang/core/ir/module.h" |
| 35 | +#include "src/tint/lang/core/ir/override.h" |
| 36 | +#include "src/tint/lang/core/ir/type/array_count.h" |
| 37 | +#include "src/tint/lang/core/ir/user_call.h" |
| 38 | +#include "src/tint/lang/core/ir/var.h" |
| 39 | +#include "src/tint/lang/core/type/array.h" |
| 40 | +#include "src/tint/lang/core/type/pointer.h" |
| 41 | +#include "src/tint/utils/containers/hashmap.h" |
| 42 | +#include "src/tint/utils/containers/unique_vector.h" |
| 43 | +#include "src/tint/utils/rtti/switch.h" |
| 44 | + |
| 45 | +// Forward declarations. |
| 46 | +namespace tint::core::ir { |
| 47 | +class Block; |
| 48 | +class Function; |
| 49 | +} // namespace tint::core::ir |
| 50 | + |
| 51 | +/// Utility that helps guarantee the same const-ness is applied to both types. |
| 52 | +template <class Src, class Dst> |
| 53 | +using TranscribeConst = std::conditional_t<std::is_const<Src>{}, std::add_const_t<Dst>, Dst>; |
| 54 | + |
| 55 | +namespace tint::core::ir { |
| 56 | + |
| 57 | +/// ReferencedModuleDecls is a helper to determine the set of module-scope declarations that are |
| 58 | +/// transitively referenced by functions in a module. |
| 59 | +/// |
| 60 | +/// References are determined lazily and cached for future requests. |
| 61 | +/// |
| 62 | +/// Note: |
| 63 | +/// The template param M is used to ensure that inputs and outputs of this class have the same |
| 64 | +/// const-ness. If 'Module' is supplied then the internal operations and output will not be |
| 65 | +/// const, which is needed for transforms. Whereas if the param is 'const Module' the internals |
| 66 | +/// and outputs will be const, which is needed for the IR validator. |
| 67 | +/// Note: |
| 68 | +/// Changes to the module can invalidate the cached data. This is intended to be created by |
| 69 | +/// operations that need this information, and discarded when they complete. Tracking this |
| 70 | +/// information inside the IR module would add overhead any time an instruction is added or |
| 71 | +/// removed from the module. Since only a few operations need this information, it is expected |
| 72 | +/// to be more efficient to generate it on demand. |
| 73 | +template <typename M> |
| 74 | +class ReferencedModuleDecls { |
| 75 | + // Replace this with concepts when C++20 is available |
| 76 | + static_assert(std::is_same<std::remove_cv_t<M>, Module>()); |
| 77 | + |
| 78 | + public: |
| 79 | + /// Short form aliases for types that have the same constant-ness as M. |
| 80 | + /// (The single use types are not aliased) |
| 81 | + using BlockT = TranscribeConst<M, Block>; |
| 82 | + using DeclT = TranscribeConst<M, Instruction>; |
| 83 | + using FunctionT = TranscribeConst<M, Function>; |
| 84 | + |
| 85 | + /// A set of a declarations referenced by a function (in declaration order). |
| 86 | + using DeclSet = UniqueVector<DeclT*, 16>; |
| 87 | + |
| 88 | + /// Constructor. |
| 89 | + /// @param ir the module |
| 90 | + explicit ReferencedModuleDecls(M& ir) { |
| 91 | + // Loop over module-scope declarations, recording the blocks that they are referenced from. |
| 92 | + BlockT* root_block = ir.root_block; |
| 93 | + for (auto* inst : *root_block) { |
| 94 | + if (!inst || !inst->Result(0)) { |
| 95 | + continue; |
| 96 | + } |
| 97 | + |
| 98 | + inst->Result(0)->ForEachUseUnsorted([&](const Usage& use) { |
| 99 | + auto& decls = block_to_direct_decls_.GetOrAddZero(use.instruction->Block()); |
| 100 | + |
| 101 | + // If this is an override we need to add the initializer to used instructions |
| 102 | + if (inst->template Is<core::ir::Override>()) { |
| 103 | + AddToBlock(decls, inst); |
| 104 | + } else { |
| 105 | + decls.Add(inst); |
| 106 | + } |
| 107 | + }); |
| 108 | + |
| 109 | + // If the instruction is a `var<workgroup>` we have to check the type. If the type is an |
| 110 | + // array with a `ValueArrayCount` then we need to check the count. If it's not |
| 111 | + // `Constant` we need to add the instruction and any referenced instructions to the used |
| 112 | + // set. |
| 113 | + auto* var = inst->template As<core::ir::Var>(); |
| 114 | + if (!var) { |
| 115 | + continue; |
| 116 | + } |
| 117 | + auto* ptr = var->Result(0)->Type()->template As<core::type::Pointer>(); |
| 118 | + TINT_ASSERT(ptr); |
| 119 | + |
| 120 | + if (ptr->AddressSpace() != core::AddressSpace::kWorkgroup) { |
| 121 | + continue; |
| 122 | + } |
| 123 | + auto* ary = ptr->UnwrapPtr()->template As<core::type::Array>(); |
| 124 | + if (!ary) { |
| 125 | + continue; |
| 126 | + } |
| 127 | + auto* cnt = ary->Count()->template As<core::ir::type::ValueArrayCount>(); |
| 128 | + if (!cnt || cnt->value->template Is<core::ir::Constant>()) { |
| 129 | + continue; |
| 130 | + } |
| 131 | + |
| 132 | + auto* cnt_inst = cnt->value->template As<core::ir::InstructionResult>(); |
| 133 | + TINT_ASSERT(cnt_inst); |
| 134 | + |
| 135 | + // The usage of the var is as a `let` initializer. The array count needs to |
| 136 | + // propagate to the `let` block. |
| 137 | + var->Result(0)->ForEachUseUnsorted([&](const Usage& use) { |
| 138 | + AddToBlock(block_to_direct_decls_.GetOrAddZero(use.instruction->Block()), |
| 139 | + cnt_inst->Instruction()); |
| 140 | + }); |
| 141 | + } |
| 142 | + } |
| 143 | + |
| 144 | + /// Get the set of transitively referenced module-scope declarations for a function. |
| 145 | + /// @param func the function |
| 146 | + /// @returns the set of transitively reference module-scope declarations |
| 147 | + DeclSet& TransitiveReferences(FunctionT* func) { |
| 148 | + return transitive_references_.GetOrAdd(func, [&] { |
| 149 | + DeclSet decls; |
| 150 | + GetTransitiveReferences(func ? func->Block() : nullptr, decls); |
| 151 | + |
| 152 | + // For a compute entry point, we need to check if any of the workgroup sizes are built |
| 153 | + // on overrides. |
| 154 | + if (func && func->Stage() == core::ir::Function::PipelineStage::kCompute) { |
| 155 | + TINT_ASSERT(func->WorkgroupSize().has_value()); |
| 156 | + |
| 157 | + const auto workgroup_size = func->WorkgroupSize(); |
| 158 | + for (auto wg_size : *workgroup_size) { |
| 159 | + if (wg_size->template Is<core::ir::Constant>()) { |
| 160 | + continue; |
| 161 | + } |
| 162 | + |
| 163 | + // Workgroup size is based on instructions, walk up the chain adding those |
| 164 | + // instructions to the `decls` list. |
| 165 | + auto* inst = wg_size->template As<core::ir::InstructionResult>(); |
| 166 | + TINT_ASSERT(inst); |
| 167 | + |
| 168 | + AddToBlock(decls, inst->Instruction()); |
| 169 | + } |
| 170 | + } |
| 171 | + return decls; |
| 172 | + }); |
| 173 | + } |
| 174 | + |
| 175 | + void AddToBlock(DeclSet& decls, core::ir::Instruction* inst) { |
| 176 | + Vector<DeclT*, 4> worklist; |
| 177 | + worklist.Push(inst); |
| 178 | + |
| 179 | + while (!worklist.IsEmpty()) { |
| 180 | + auto* wl_inst = worklist.Pop(); |
| 181 | + if (decls.Add(wl_inst)) { |
| 182 | + for (auto* operand : wl_inst->Operands()) { |
| 183 | + if (!operand) { |
| 184 | + continue; |
| 185 | + } |
| 186 | + auto* res = operand->template As<core::ir::InstructionResult>(); |
| 187 | + if (!res) { |
| 188 | + continue; |
| 189 | + } |
| 190 | + worklist.Push(res->Instruction()); |
| 191 | + } |
| 192 | + } |
| 193 | + |
| 194 | + for (auto* operand : wl_inst->Operands()) { |
| 195 | + if (!operand) { |
| 196 | + continue; |
| 197 | + } |
| 198 | + auto* res = operand->template As<core::ir::InstructionResult>(); |
| 199 | + if (!res) { |
| 200 | + continue; |
| 201 | + } |
| 202 | + worklist.Push(res->Instruction()); |
| 203 | + } |
| 204 | + } |
| 205 | + } |
| 206 | + |
| 207 | + private: |
| 208 | + /// A map from blocks to their directly referenced declarations. |
| 209 | + Hashmap<BlockT*, DeclSet, 64> block_to_direct_decls_{}; |
| 210 | + |
| 211 | + /// A map from functions to their transitively referenced declarations. |
| 212 | + Hashmap<FunctionT*, DeclSet, 8> transitive_references_; |
| 213 | + |
| 214 | + /// Get the set of transitively referenced module-scope declarations for a block. |
| 215 | + /// @param block the block |
| 216 | + /// @param decls the set of transitively reference module-scope declarations to populate |
| 217 | + void GetTransitiveReferences(BlockT* block, DeclSet& decls) { |
| 218 | + if (!block) { |
| 219 | + return; |
| 220 | + } |
| 221 | + |
| 222 | + // Add directly referenced declarations. |
| 223 | + if (auto itr = block_to_direct_decls_.Get(block)) { |
| 224 | + for (auto& decl : *itr) { |
| 225 | + decls.Add(decl); |
| 226 | + } |
| 227 | + } |
| 228 | + |
| 229 | + // Loop over instructions in the block to find indirectly referenced vars. |
| 230 | + for (auto* inst : *block) { |
| 231 | + tint::Switch( |
| 232 | + inst, |
| 233 | + [&](TranscribeConst<M, UserCall>* call) { |
| 234 | + // Get declarations referenced by a function called from this block. |
| 235 | + const auto& callee_decls = TransitiveReferences(call->Target()); |
| 236 | + for (auto* decl : callee_decls) { |
| 237 | + decls.Add(decl); |
| 238 | + } |
| 239 | + }, |
| 240 | + [&](TranscribeConst<M, ControlInstruction>* ctrl) { |
| 241 | + // Recurse into control instructions and gather their referenced declarations. |
| 242 | + ctrl->ForeachBlock([&](BlockT* blk) { GetTransitiveReferences(blk, decls); }); |
| 243 | + }); |
| 244 | + } |
| 245 | + } |
| 246 | +}; |
| 247 | + |
| 248 | +} // namespace tint::core::ir |
| 249 | + |
| 250 | +#endif // SRC_TINT_LANG_CORE_IR_REFERENCED_MODULE_DECLS_H_ |
0 commit comments