| // Copyright 2022 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/lang/core/ir/disassembler.h" |
| |
| #include "src//tint/lang/core/ir/unary.h" |
| #include "src/tint/lang/core/constant/composite.h" |
| #include "src/tint/lang/core/constant/scalar.h" |
| #include "src/tint/lang/core/constant/splat.h" |
| #include "src/tint/lang/core/fluent_types.h" |
| #include "src/tint/lang/core/ir/access.h" |
| #include "src/tint/lang/core/ir/binary.h" |
| #include "src/tint/lang/core/ir/bitcast.h" |
| #include "src/tint/lang/core/ir/block.h" |
| #include "src/tint/lang/core/ir/block_param.h" |
| #include "src/tint/lang/core/ir/break_if.h" |
| #include "src/tint/lang/core/ir/construct.h" |
| #include "src/tint/lang/core/ir/continue.h" |
| #include "src/tint/lang/core/ir/convert.h" |
| #include "src/tint/lang/core/ir/core_builtin_call.h" |
| #include "src/tint/lang/core/ir/discard.h" |
| #include "src/tint/lang/core/ir/exit_if.h" |
| #include "src/tint/lang/core/ir/exit_loop.h" |
| #include "src/tint/lang/core/ir/exit_switch.h" |
| #include "src/tint/lang/core/ir/if.h" |
| #include "src/tint/lang/core/ir/instruction_result.h" |
| #include "src/tint/lang/core/ir/let.h" |
| #include "src/tint/lang/core/ir/load.h" |
| #include "src/tint/lang/core/ir/load_vector_element.h" |
| #include "src/tint/lang/core/ir/loop.h" |
| #include "src/tint/lang/core/ir/multi_in_block.h" |
| #include "src/tint/lang/core/ir/next_iteration.h" |
| #include "src/tint/lang/core/ir/return.h" |
| #include "src/tint/lang/core/ir/store.h" |
| #include "src/tint/lang/core/ir/store_vector_element.h" |
| #include "src/tint/lang/core/ir/switch.h" |
| #include "src/tint/lang/core/ir/swizzle.h" |
| #include "src/tint/lang/core/ir/terminate_invocation.h" |
| #include "src/tint/lang/core/ir/unreachable.h" |
| #include "src/tint/lang/core/ir/user_call.h" |
| #include "src/tint/lang/core/ir/var.h" |
| #include "src/tint/lang/core/type/struct.h" |
| #include "src/tint/lang/core/type/type.h" |
| #include "src/tint/utils/ice/ice.h" |
| #include "src/tint/utils/macros/scoped_assignment.h" |
| #include "src/tint/utils/rtti/switch.h" |
| #include "src/tint/utils/text/string.h" |
| |
| using namespace tint::core::fluent_types; // NOLINT |
| |
| namespace tint::core::ir { |
| namespace { |
| |
| class ScopedIndent { |
| public: |
| explicit ScopedIndent(uint32_t& indent) : indent_(indent) { indent_ += 2; } |
| |
| ~ScopedIndent() { indent_ -= 2; } |
| |
| private: |
| uint32_t& indent_; |
| }; |
| |
| } // namespace |
| |
| Disassembler::Disassembler(Module& mod) : mod_(mod) {} |
| |
| Disassembler::~Disassembler() = default; |
| |
| StringStream& Disassembler::Indent() { |
| for (uint32_t i = 0; i < indent_size_; i++) { |
| out_ << " "; |
| } |
| return out_; |
| } |
| |
| void Disassembler::EmitLine() { |
| out_ << std::endl; |
| current_output_line_ += 1; |
| current_output_start_pos_ = out_.tellp(); |
| } |
| |
| size_t Disassembler::IdOf(Block* node) { |
| TINT_ASSERT(node); |
| return block_ids_.GetOrCreate(node, [&] { return block_ids_.Count(); }); |
| } |
| |
| std::string Disassembler::IdOf(Value* value) { |
| TINT_ASSERT(value); |
| return value_ids_.GetOrCreate(value, [&] { |
| if (auto sym = mod_.NameOf(value)) { |
| if (ids_.Add(sym.Name())) { |
| return sym.Name(); |
| } |
| auto prefix = sym.Name() + "_"; |
| for (size_t i = 1;; i++) { |
| auto name = prefix + std::to_string(i); |
| if (ids_.Add(name)) { |
| return name; |
| } |
| } |
| } |
| return std::to_string(value_ids_.Count()); |
| }); |
| } |
| |
| std::string Disassembler::NameOf(If* inst) { |
| if (!inst) { |
| return "undef"; |
| } |
| |
| return if_names_.GetOrCreate(inst, [&] { return "if_" + std::to_string(if_names_.Count()); }); |
| } |
| |
| std::string Disassembler::NameOf(Loop* inst) { |
| if (!inst) { |
| return "undef"; |
| } |
| |
| return loop_names_.GetOrCreate(inst, |
| [&] { return "loop_" + std::to_string(loop_names_.Count()); }); |
| } |
| |
| std::string Disassembler::NameOf(Switch* inst) { |
| if (!inst) { |
| return "undef"; |
| } |
| |
| return switch_names_.GetOrCreate( |
| inst, [&] { return "switch_" + std::to_string(switch_names_.Count()); }); |
| } |
| |
| Source::Location Disassembler::MakeCurrentLocation() { |
| return Source::Location{current_output_line_, out_.tellp() - current_output_start_pos_ + 1}; |
| } |
| |
| std::string Disassembler::Disassemble() { |
| for (auto* ty : mod_.Types()) { |
| if (auto* str = ty->As<core::type::Struct>()) { |
| EmitStructDecl(str); |
| } |
| } |
| |
| if (!mod_.root_block->IsEmpty()) { |
| EmitBlock(mod_.root_block, "root"); |
| EmitLine(); |
| } |
| |
| for (auto* func : mod_.functions) { |
| EmitFunction(func); |
| } |
| return out_.str(); |
| } |
| |
| void Disassembler::EmitBlock(Block* blk, std::string_view comment /* = "" */) { |
| Indent(); |
| |
| SourceMarker sm(this); |
| out_ << "%b" << IdOf(blk) << " = block"; |
| if (auto* merge = blk->As<MultiInBlock>()) { |
| if (!merge->Params().IsEmpty()) { |
| out_ << " ("; |
| EmitValueList(merge->Params().Slice()); |
| out_ << ")"; |
| } |
| } |
| sm.Store(blk); |
| |
| out_ << " {"; |
| if (!comment.empty()) { |
| out_ << " # " << comment; |
| } |
| |
| EmitLine(); |
| { |
| ScopedIndent si(indent_size_); |
| for (auto* inst : *blk) { |
| Indent(); |
| EmitInstruction(inst); |
| } |
| } |
| Indent() << "}"; |
| |
| EmitLine(); |
| } |
| |
| void Disassembler::EmitBindingPoint(BindingPoint p) { |
| out_ << "@binding_point(" << p.group << ", " << p.binding << ")"; |
| } |
| |
| void Disassembler::EmitLocation(Location loc) { |
| out_ << "@location(" << loc.value << ")"; |
| if (loc.interpolation.has_value()) { |
| out_ << ", @interpolate("; |
| out_ << loc.interpolation->type; |
| if (loc.interpolation->sampling != core::InterpolationSampling::kUndefined) { |
| out_ << ", "; |
| out_ << loc.interpolation->sampling; |
| } |
| out_ << ")"; |
| } |
| } |
| |
| void Disassembler::EmitParamAttributes(FunctionParam* p) { |
| if (!p->Invariant() && !p->Location().has_value() && !p->BindingPoint().has_value() && |
| !p->Builtin().has_value()) { |
| return; |
| } |
| |
| out_ << " ["; |
| |
| bool need_comma = false; |
| auto comma = [&] { |
| if (need_comma) { |
| out_ << ", "; |
| } |
| }; |
| |
| if (p->Invariant()) { |
| comma(); |
| out_ << "@invariant"; |
| need_comma = true; |
| } |
| if (p->Location().has_value()) { |
| EmitLocation(p->Location().value()); |
| need_comma = true; |
| } |
| if (p->BindingPoint().has_value()) { |
| comma(); |
| EmitBindingPoint(p->BindingPoint().value()); |
| need_comma = true; |
| } |
| if (p->Builtin().has_value()) { |
| comma(); |
| out_ << "@" << p->Builtin().value(); |
| need_comma = true; |
| } |
| out_ << "]"; |
| } |
| |
| void Disassembler::EmitReturnAttributes(Function* func) { |
| if (!func->ReturnInvariant() && !func->ReturnLocation().has_value() && |
| !func->ReturnBuiltin().has_value()) { |
| return; |
| } |
| |
| out_ << " ["; |
| |
| bool need_comma = false; |
| auto comma = [&] { |
| if (need_comma) { |
| out_ << ", "; |
| } |
| }; |
| if (func->ReturnInvariant()) { |
| comma(); |
| out_ << "@invariant"; |
| need_comma = true; |
| } |
| if (func->ReturnLocation().has_value()) { |
| comma(); |
| EmitLocation(func->ReturnLocation().value()); |
| need_comma = true; |
| } |
| if (func->ReturnBuiltin().has_value()) { |
| comma(); |
| out_ << "@" << func->ReturnBuiltin().value(); |
| need_comma = true; |
| } |
| out_ << "]"; |
| } |
| |
| void Disassembler::EmitFunction(Function* func) { |
| in_function_ = true; |
| |
| std::string fn_id = IdOf(func); |
| Indent() << "%" << fn_id << " ="; |
| |
| if (func->Stage() != Function::PipelineStage::kUndefined) { |
| out_ << " @" << func->Stage(); |
| } |
| if (func->WorkgroupSize()) { |
| auto arr = func->WorkgroupSize().value(); |
| out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] << ")"; |
| } |
| |
| out_ << " func("; |
| |
| for (auto* p : func->Params()) { |
| if (p != func->Params().Front()) { |
| out_ << ", "; |
| } |
| out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName(); |
| |
| EmitParamAttributes(p); |
| } |
| out_ << "):" << func->ReturnType()->FriendlyName(); |
| |
| EmitReturnAttributes(func); |
| |
| out_ << " -> %b" << IdOf(func->Block()) << " {"; |
| |
| { // Add a comment if the function IDs or parameter IDs doesn't match their name |
| Vector<std::string, 4> names; |
| if (auto name = mod_.NameOf(func); name.IsValid()) { |
| if (name.NameView() != fn_id) { |
| names.Push("%" + std::string(fn_id) + ": '" + name.Name() + "'"); |
| } |
| } |
| for (auto* p : func->Params()) { |
| if (auto name = mod_.NameOf(p); name.IsValid()) { |
| auto id = IdOf(p); |
| if (name.NameView() != id) { |
| names.Push("%" + std::string(id) + ": '" + name.Name() + "'"); |
| } |
| } |
| } |
| if (!names.IsEmpty()) { |
| out_ << " # " << tint::Join(names, ", "); |
| } |
| } |
| |
| EmitLine(); |
| |
| { |
| ScopedIndent si(indent_size_); |
| EmitBlock(func->Block()); |
| } |
| Indent() << "}"; |
| EmitLine(); |
| } |
| |
| void Disassembler::EmitValueWithType(Instruction* val) { |
| SourceMarker sm(this); |
| if (val->Result()) { |
| EmitValueWithType(val->Result()); |
| } else { |
| out_ << "undef"; |
| } |
| sm.StoreResult(Usage{val, 0}); |
| } |
| |
| void Disassembler::EmitValueWithType(Value* val) { |
| if (!val) { |
| out_ << "undef"; |
| return; |
| } |
| |
| EmitValue(val); |
| out_ << ":" << val->Type()->FriendlyName(); |
| } |
| |
| void Disassembler::EmitValue(Value* val) { |
| tint::Switch( |
| val, |
| [&](ir::Constant* constant) { |
| std::function<void(const core::constant::Value*)> emit = |
| [&](const core::constant::Value* c) { |
| tint::Switch( |
| c, |
| [&](const core::constant::Scalar<AFloat>* scalar) { |
| out_ << scalar->ValueAs<AFloat>().value; |
| }, |
| [&](const core::constant::Scalar<AInt>* scalar) { |
| out_ << scalar->ValueAs<AInt>().value; |
| }, |
| [&](const core::constant::Scalar<i32>* scalar) { |
| out_ << scalar->ValueAs<i32>().value << "i"; |
| }, |
| [&](const core::constant::Scalar<u32>* scalar) { |
| out_ << scalar->ValueAs<u32>().value << "u"; |
| }, |
| [&](const core::constant::Scalar<f32>* scalar) { |
| out_ << scalar->ValueAs<f32>().value << "f"; |
| }, |
| [&](const core::constant::Scalar<f16>* scalar) { |
| out_ << scalar->ValueAs<f16>().value << "h"; |
| }, |
| [&](const core::constant::Scalar<bool>* scalar) { |
| out_ << (scalar->ValueAs<bool>() ? "true" : "false"); |
| }, |
| [&](const core::constant::Splat* splat) { |
| out_ << splat->Type()->FriendlyName() << "("; |
| emit(splat->Index(0)); |
| out_ << ")"; |
| }, |
| [&](const core::constant::Composite* composite) { |
| out_ << composite->Type()->FriendlyName() << "("; |
| bool need_comma = false; |
| for (const auto* elem : composite->elements) { |
| if (need_comma) { |
| out_ << ", "; |
| } |
| emit(elem); |
| need_comma = true; |
| } |
| out_ << ")"; |
| }); |
| }; |
| emit(constant->Value()); |
| }, |
| [&](ir::InstructionResult* rv) { out_ << "%" << IdOf(rv); }, |
| [&](ir::BlockParam* p) { out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName(); }, |
| [&](ir::FunctionParam* p) { out_ << "%" << IdOf(p); }, |
| [&](ir::Function* f) { out_ << "%" << IdOf(f); }, |
| [&](Default) { |
| if (val == nullptr) { |
| out_ << "undef"; |
| } else { |
| out_ << "Unknown value: " << val->TypeInfo().name; |
| } |
| }); |
| } |
| |
| void Disassembler::EmitInstructionName(Instruction* inst) { |
| SourceMarker sm(this); |
| out_ << inst->FriendlyName(); |
| sm.Store(inst); |
| } |
| |
| void Disassembler::EmitInstruction(Instruction* inst) { |
| TINT_DEFER(EmitLine()); |
| |
| if (!inst->Alive()) { |
| SourceMarker sm(this); |
| out_ << "<destroyed " << inst->TypeInfo().name << " " << tint::ToString(inst) << ">"; |
| sm.Store(inst); |
| return; |
| } |
| tint::Switch( |
| inst, // |
| [&](Switch* s) { EmitSwitch(s); }, // |
| [&](If* i) { EmitIf(i); }, // |
| [&](Loop* l) { EmitLoop(l); }, // |
| [&](Binary* b) { EmitBinary(b); }, // |
| [&](Unary* u) { EmitUnary(u); }, // |
| [&](Discard* d) { EmitInstructionName(d); }, |
| [&](Store* s) { |
| EmitInstructionName(s); |
| out_ << " "; |
| EmitOperand(s, Store::kToOperandOffset); |
| out_ << ", "; |
| EmitOperand(s, Store::kFromOperandOffset); |
| }, |
| [&](StoreVectorElement* s) { |
| EmitInstructionName(s); |
| out_ << " "; |
| EmitOperandList(s); |
| }, |
| [&](UserCall* uc) { |
| EmitValueWithType(uc); |
| out_ << " = "; |
| EmitInstructionName(uc); |
| out_ << " "; |
| EmitOperand(uc, UserCall::kFunctionOperandOffset); |
| if (!uc->Args().IsEmpty()) { |
| out_ << ", "; |
| } |
| EmitOperandList(uc, UserCall::kArgsOperandOffset); |
| }, |
| [&](Var* v) { |
| EmitValueWithType(v); |
| out_ << " = "; |
| EmitInstructionName(v); |
| if (v->Initializer()) { |
| out_ << ", "; |
| EmitOperand(v, Var::kInitializerOperandOffset); |
| } |
| if (v->BindingPoint().has_value()) { |
| out_ << " "; |
| EmitBindingPoint(v->BindingPoint().value()); |
| } |
| if (v->Attributes().invariant) { |
| out_ << " @invariant"; |
| } |
| if (v->Attributes().location.has_value()) { |
| out_ << " @location(" << v->Attributes().location.value() << ")"; |
| } |
| if (v->Attributes().index.has_value()) { |
| out_ << " @index(" << v->Attributes().index.value() << ")"; |
| } |
| if (v->Attributes().interpolation.has_value()) { |
| auto& interp = v->Attributes().interpolation.value(); |
| out_ << " @interpolate(" << interp.type; |
| if (interp.sampling != core::InterpolationSampling::kUndefined) { |
| out_ << ", " << interp.sampling; |
| } |
| out_ << ")"; |
| } |
| if (v->Attributes().builtin.has_value()) { |
| out_ << " @builtin(" << v->Attributes().builtin.value() << ")"; |
| } |
| }, |
| [&](Swizzle* s) { |
| EmitValueWithType(s); |
| out_ << " = "; |
| EmitInstructionName(s); |
| out_ << " "; |
| EmitValue(s->Object()); |
| out_ << ", "; |
| for (auto idx : s->Indices()) { |
| switch (idx) { |
| case 0: |
| out_ << "x"; |
| break; |
| case 1: |
| out_ << "y"; |
| break; |
| case 2: |
| out_ << "z"; |
| break; |
| case 3: |
| out_ << "w"; |
| break; |
| } |
| } |
| }, |
| [&](Terminator* b) { EmitTerminator(b); }, |
| [&](Default) { |
| EmitValueWithType(inst); |
| out_ << " = "; |
| EmitInstructionName(inst); |
| if (!inst->Operands().IsEmpty()) { |
| out_ << " "; |
| EmitOperandList(inst); |
| } |
| }); |
| |
| { // Add a comment if the result IDs don't match their names |
| Vector<std::string, 4> names; |
| for (auto* result : inst->Results()) { |
| if (result) { |
| if (auto name = mod_.NameOf(result); name.IsValid()) { |
| auto id = IdOf(result); |
| if (name.NameView() != id) { |
| names.Push("%" + std::string(id) + ": '" + name.Name() + "'"); |
| } |
| } |
| } |
| } |
| if (!names.IsEmpty()) { |
| out_ << " # " << tint::Join(names, ", "); |
| } |
| } |
| } |
| |
| void Disassembler::EmitOperand(Instruction* inst, size_t index) { |
| SourceMarker condMarker(this); |
| EmitValue(inst->Operands()[index]); |
| condMarker.Store(Usage{inst, static_cast<uint32_t>(index)}); |
| } |
| |
| void Disassembler::EmitOperandList(Instruction* inst, size_t start_index /* = 0 */) { |
| for (size_t i = start_index, n = inst->Operands().Length(); i < n; i++) { |
| if (i != start_index) { |
| out_ << ", "; |
| } |
| EmitOperand(inst, i); |
| } |
| } |
| |
| void Disassembler::EmitIf(If* if_) { |
| SourceMarker sm(this); |
| if (if_->HasResults()) { |
| auto res = if_->Results(); |
| for (size_t i = 0; i < res.Length(); ++i) { |
| if (i > 0) { |
| out_ << ", "; |
| } |
| SourceMarker rs(this); |
| EmitValueWithType(res[i]); |
| rs.StoreResult(Usage{if_, i}); |
| } |
| out_ << " = "; |
| } |
| out_ << "if "; |
| EmitOperand(if_, If::kConditionOperandOffset); |
| |
| bool has_false = !if_->False()->IsEmpty(); |
| |
| out_ << " [t: %b" << IdOf(if_->True()); |
| if (has_false) { |
| out_ << ", f: %b" << IdOf(if_->False()); |
| } |
| out_ << "]"; |
| sm.Store(if_); |
| |
| out_ << " { # " << NameOf(if_); |
| EmitLine(); |
| |
| // True block is assumed to have instructions |
| { |
| ScopedIndent si(indent_size_); |
| EmitBlock(if_->True(), "true"); |
| } |
| |
| if (has_false) { |
| ScopedIndent si(indent_size_); |
| EmitBlock(if_->False(), "false"); |
| } else if (if_->HasResults()) { |
| ScopedIndent si(indent_size_); |
| Indent(); |
| out_ << "# implicit false block: exit_if undef"; |
| for (size_t v = 1; v < if_->Results().Length(); v++) { |
| out_ << ", undef"; |
| } |
| EmitLine(); |
| } |
| |
| Indent(); |
| out_ << "}"; |
| } |
| |
| void Disassembler::EmitLoop(Loop* l) { |
| Vector<std::string, 3> parts; |
| if (!l->Initializer()->IsEmpty()) { |
| parts.Push("i: %b" + std::to_string(IdOf(l->Initializer()))); |
| } |
| parts.Push("b: %b" + std::to_string(IdOf(l->Body()))); |
| |
| if (!l->Continuing()->IsEmpty()) { |
| parts.Push("c: %b" + std::to_string(IdOf(l->Continuing()))); |
| } |
| SourceMarker sm(this); |
| if (l->HasResults()) { |
| auto res = l->Results(); |
| for (size_t i = 0; i < res.Length(); ++i) { |
| if (i > 0) { |
| out_ << ", "; |
| } |
| SourceMarker rs(this); |
| EmitValueWithType(res[i]); |
| rs.StoreResult(Usage{l, i}); |
| } |
| out_ << " = "; |
| } |
| out_ << "loop [" << tint::Join(parts, ", ") << "]"; |
| sm.Store(l); |
| |
| out_ << " { # " << NameOf(l); |
| EmitLine(); |
| |
| if (!l->Initializer()->IsEmpty()) { |
| ScopedIndent si(indent_size_); |
| EmitBlock(l->Initializer(), "initializer"); |
| } |
| |
| // Loop is assumed to always have a body |
| { |
| ScopedIndent si(indent_size_); |
| EmitBlock(l->Body(), "body"); |
| } |
| |
| if (!l->Continuing()->IsEmpty()) { |
| ScopedIndent si(indent_size_); |
| EmitBlock(l->Continuing(), "continuing"); |
| } |
| |
| Indent(); |
| out_ << "}"; |
| } |
| |
| void Disassembler::EmitSwitch(Switch* s) { |
| SourceMarker sm(this); |
| if (s->HasResults()) { |
| auto res = s->Results(); |
| for (size_t i = 0; i < res.Length(); ++i) { |
| if (i > 0) { |
| out_ << ", "; |
| } |
| SourceMarker rs(this); |
| EmitValueWithType(res[i]); |
| rs.StoreResult(Usage{s, i}); |
| } |
| out_ << " = "; |
| } |
| out_ << "switch "; |
| EmitValue(s->Condition()); |
| out_ << " ["; |
| for (auto& c : s->Cases()) { |
| if (&c != &s->Cases().Front()) { |
| out_ << ", "; |
| } |
| out_ << "c: ("; |
| for (auto& selector : c.selectors) { |
| if (&selector != &c.selectors.Front()) { |
| out_ << " "; |
| } |
| |
| if (selector.IsDefault()) { |
| out_ << "default"; |
| } else { |
| EmitValue(selector.val); |
| } |
| } |
| out_ << ", %b" << IdOf(c.Block()) << ")"; |
| } |
| out_ << "]"; |
| sm.Store(s); |
| |
| out_ << " { # " << NameOf(s); |
| EmitLine(); |
| |
| for (auto& c : s->Cases()) { |
| ScopedIndent si(indent_size_); |
| EmitBlock(c.Block(), "case"); |
| } |
| |
| Indent(); |
| out_ << "}"; |
| } |
| |
| void Disassembler::EmitTerminator(Terminator* b) { |
| SourceMarker sm(this); |
| size_t args_offset = 0; |
| tint::Switch( |
| b, |
| [&](ir::Return*) { |
| out_ << "ret"; |
| args_offset = ir::Return::kArgOperandOffset; |
| }, |
| [&](ir::Continue* cont) { |
| out_ << "continue %b" << IdOf(cont->Loop()->Continuing()); |
| args_offset = ir::Continue::kArgsOperandOffset; |
| }, |
| [&](ir::ExitIf*) { |
| out_ << "exit_if"; |
| args_offset = ir::ExitIf::kArgsOperandOffset; |
| }, |
| [&](ir::ExitSwitch*) { |
| out_ << "exit_switch"; |
| args_offset = ir::ExitSwitch::kArgsOperandOffset; |
| }, |
| [&](ir::ExitLoop*) { |
| out_ << "exit_loop"; |
| args_offset = ir::ExitLoop::kArgsOperandOffset; |
| }, |
| [&](ir::NextIteration* ni) { |
| out_ << "next_iteration %b" << IdOf(ni->Loop()->Body()); |
| args_offset = ir::NextIteration::kArgsOperandOffset; |
| }, |
| [&](ir::Unreachable*) { out_ << "unreachable"; }, |
| [&](ir::BreakIf* bi) { |
| out_ << "break_if "; |
| EmitValue(bi->Condition()); |
| out_ << " %b" << IdOf(bi->Loop()->Body()); |
| args_offset = ir::BreakIf::kArgsOperandOffset; |
| }, |
| [&](ir::TerminateInvocation*) { out_ << "terminate_invocation"; }, |
| [&](Default) { out_ << "unknown terminator " << b->TypeInfo().name; }); |
| |
| if (!b->Args().IsEmpty()) { |
| out_ << " "; |
| EmitOperandList(b, args_offset); |
| } |
| sm.Store(b); |
| |
| tint::Switch( |
| b, // |
| [&](ir::ExitIf* e) { out_ << " # " << NameOf(e->If()); }, // |
| [&](ir::ExitSwitch* e) { out_ << " # " << NameOf(e->Switch()); }, // |
| [&](ir::ExitLoop* e) { out_ << " # " << NameOf(e->Loop()); } // |
| ); |
| } |
| |
| void Disassembler::EmitValueList(tint::Slice<Value* const> values) { |
| for (size_t i = 0, n = values.Length(); i < n; i++) { |
| if (i > 0) { |
| out_ << ", "; |
| } |
| EmitValue(values[i]); |
| } |
| } |
| |
| void Disassembler::EmitBinary(Binary* b) { |
| SourceMarker sm(this); |
| EmitValueWithType(b); |
| out_ << " = "; |
| switch (b->Kind()) { |
| case Binary::Kind::kAdd: |
| out_ << "add"; |
| break; |
| case Binary::Kind::kSubtract: |
| out_ << "sub"; |
| break; |
| case Binary::Kind::kMultiply: |
| out_ << "mul"; |
| break; |
| case Binary::Kind::kDivide: |
| out_ << "div"; |
| break; |
| case Binary::Kind::kModulo: |
| out_ << "mod"; |
| break; |
| case Binary::Kind::kAnd: |
| out_ << "and"; |
| break; |
| case Binary::Kind::kOr: |
| out_ << "or"; |
| break; |
| case Binary::Kind::kXor: |
| out_ << "xor"; |
| break; |
| case Binary::Kind::kEqual: |
| out_ << "eq"; |
| break; |
| case Binary::Kind::kNotEqual: |
| out_ << "neq"; |
| break; |
| case Binary::Kind::kLessThan: |
| out_ << "lt"; |
| break; |
| case Binary::Kind::kGreaterThan: |
| out_ << "gt"; |
| break; |
| case Binary::Kind::kLessThanEqual: |
| out_ << "lte"; |
| break; |
| case Binary::Kind::kGreaterThanEqual: |
| out_ << "gte"; |
| break; |
| case Binary::Kind::kShiftLeft: |
| out_ << "shiftl"; |
| break; |
| case Binary::Kind::kShiftRight: |
| out_ << "shiftr"; |
| break; |
| } |
| out_ << " "; |
| EmitOperandList(b); |
| |
| sm.Store(b); |
| } |
| |
| void Disassembler::EmitUnary(Unary* u) { |
| SourceMarker sm(this); |
| EmitValueWithType(u); |
| out_ << " = "; |
| switch (u->Kind()) { |
| case Unary::Kind::kComplement: |
| out_ << "complement"; |
| break; |
| case Unary::Kind::kNegation: |
| out_ << "negation"; |
| break; |
| } |
| out_ << " "; |
| EmitOperandList(u); |
| |
| sm.Store(u); |
| } |
| |
| void Disassembler::EmitStructDecl(const core::type::Struct* str) { |
| out_ << str->Name().Name() << " = struct @align(" << str->Align() << ")"; |
| if (str->StructFlags().Contains(core::type::StructFlag::kBlock)) { |
| out_ << ", @block"; |
| } |
| out_ << " {"; |
| EmitLine(); |
| for (auto* member : str->Members()) { |
| out_ << " " << member->Name().Name() << ":" << member->Type()->FriendlyName(); |
| out_ << " @offset(" << member->Offset() << ")"; |
| if (member->Attributes().invariant) { |
| out_ << ", @invariant"; |
| } |
| if (member->Attributes().location.has_value()) { |
| out_ << ", @location(" << member->Attributes().location.value() << ")"; |
| } |
| if (member->Attributes().interpolation.has_value()) { |
| auto& interp = member->Attributes().interpolation.value(); |
| out_ << ", @interpolate(" << interp.type; |
| if (interp.sampling != core::InterpolationSampling::kUndefined) { |
| out_ << ", " << interp.sampling; |
| } |
| out_ << ")"; |
| } |
| if (member->Attributes().builtin.has_value()) { |
| out_ << ", @builtin(" << member->Attributes().builtin.value() << ")"; |
| } |
| EmitLine(); |
| } |
| out_ << "}"; |
| EmitLine(); |
| EmitLine(); |
| } |
| |
| } // namespace tint::core::ir |
