comaniac commented on a change in pull request #8615: URL: https://github.com/apache/tvm/pull/8615#discussion_r680561261
########## File path: python/tvm/tir/schedule/instruction.py ########## @@ -0,0 +1,166 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you 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. +"""Schedule instructions each corresponds to a schedule primitive""" +from typing import TYPE_CHECKING, Any, List, Union + +from tvm._ffi import register_object as _register_object +from tvm.runtime import Object + +from . import _ffi_api + +if TYPE_CHECKING: + from .schedule import RAND_VAR_TYPE + + INPUT_RV_TYPE = Union[RAND_VAR_TYPE, float, int, str, None] # pylint: disable=invalid-name + OUTPUT_RV_TYPE = Union[RAND_VAR_TYPE] # pylint: disable=invalid-name + ATTR_TYPE = Any +else: + INPUT_RV_TYPE = OUTPUT_RV_TYPE = ATTR_TYPE = Any + + +@_register_object("tir.InstructionKind") +class InstructionKind(Object): + """Kind of an instruction, e.g. Split, Reorder, etc. + Besides the name, every kind of instruction has its own properties, including: + 1) A boolean indicating if the instruction is pure, i.e. change nothing in the schedule state + 2) A functor that applies the instruction to a TensorIR schedule + 3) A functor that converts the instruction to a statement in python syntax + 4) A functor that serialize its attributes to JSON + 5) A functor that deserialize its attributes from JSON + + Unlike `tvm.ir.op`, `InstructionKind` doesn't support unstructured properties, + mainly because there is no such usecase yet to add any other property. + + Attributes + ---------- + name : str + The name of a kind of instructions + + Note + ---------- Review comment: ```suggestion ---- ``` ########## File path: python/tvm/tir/schedule/trace.py ########## @@ -0,0 +1,260 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you 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. +"""An execution trace of a scheduling program""" +from typing import TYPE_CHECKING, Any, Callable, Dict, List, Optional + +from tvm._ffi import register_object as _register_object +from tvm.runtime import Object + +from ...ir import Array, Map +from ...runtime import String +from ..expr import FloatImm, IntImm +from . import _ffi_api +from .instruction import ATTR_TYPE, INPUT_RV_TYPE, Instruction + +if TYPE_CHECKING: + from .schedule import ExprRV, Schedule + + +DECISION_TYPE = Any +JSON_TYPE = Any + + +def _json_from_tvm(obj): + if obj is None: + return None + if isinstance(obj, Array): + return [_json_from_tvm(i) for i in obj] + if isinstance(obj, Map): + return {_json_from_tvm(k): _json_from_tvm(v) for k, v in obj.items()} + if isinstance(obj, String): + return str(obj) + if isinstance(obj, (IntImm, FloatImm)): + return obj.value + raise TypeError("Not supported type: " + str(type(obj))) + + +@_register_object("tir.Trace") +class Trace(Object): + """An execution trace of a scheduling program. + + A trace has two parts: + 1) The instructions invoked so far + 2) The random decisions made upon those instructions, if any + + A trace can be serialized to: + 1) Roundtrippable JSON format: can be saved to file and loaded back + 2) Python syntax: allows users to copy-paste the trace to reproduce the scheduling process + + A trace can be applied to a TensorIR schedule by re-applying all its instructions possibly with + their decisions accordingly. Re-sampling is invoked if a sampling instruction doesn't have its + corresponding decision; Otherwise the existing decision will be reused accordingly. + + Attributes + ---------- + insts : List[Instruction] + The instructions invoked so far in the program execution + decisions : Dict[Instruction, DECISION_TYPE] + The random decisions made upon those instructions + """ + + insts: List[Instruction] + decisions: Dict[Instruction, DECISION_TYPE] + + def __init__( + self, + insts: List[Instruction], + decisions: Dict[Instruction, DECISION_TYPE], + ) -> None: + """Constructor + + Parameters + ---------- + insts : List[Instruction] + The instructions invoked so far in the program execution + decisions : Dict[Instruction, DECISION_TYPE] + The random decisions made upon those instructions + """ + self.__init_handle_by_constructor__( + _ffi_api.Trace, # type: ignore # pylint: disable=no-member + insts, + decisions, + ) + + def get_decision(self, inst: Instruction) -> Optional[DECISION_TYPE]: + """Retrieve the decision made on a specific instruction + + Parameters + ---------- + insts : Instruction + The instruction whose decision is to be retrieved + + Returns + ---------- + decision : Optional[DECISION_TYPE] + The corresponding decision; None if there is no decision made on the instruction + """ + return _ffi_api.TraceGetDecision(self, inst) # type: ignore # pylint: disable=no-member + + def append( + self, + inst: Instruction, + decision: Optional[DECISION_TYPE] = None, + ) -> None: + """Append a new instruction to the trace + + Parameters + ---------- + insts : Instruction + The new instruction to be appended + decision : Optional[DECISION_TYPE] = None + The random decision made on this instruction + """ + _ffi_api.TraceAppend(self, inst, decision) # type: ignore # pylint: disable=no-member + + def pop(self) -> Optional[Instruction]: + """Remove the last instruction, along with the decision made on that instruction, if any + + Returns + ---------- + popped_inst : Instruction + Returns the instruction removed; NullOpt if the trace is empty + """ + return _ffi_api.TracePop(self) # type: ignore # pylint: disable=no-member + + def apply_to_schedule( + self, + sch: "Schedule", + remove_postproc: bool, + decision_provider: Optional[ + Callable[ + [Instruction, List[INPUT_RV_TYPE], List[ATTR_TYPE], DECISION_TYPE], DECISION_TYPE + ] + ] = None, + ) -> None: + """Apply the trace to a TensorIR schedule + + Parameters + ---------- + sch : Schedule + The schedule to be applied onto + remove_postproc : bool + If postprocessing instructions are removed + decision_provider: Optional[Callable] = None + A callback that allows users to mutate decisions on the fly when applying instructions. + The signature of the callback is: + - The 1st argument: The instruction + - The 2nd argument: The input random variables + - The 3rd argument: The attributes + - The 4th argument: The decision + - Return: A new decision + """ + _ffi_api.TraceApplyToSchedule( # type: ignore # pylint: disable=no-member + self, + sch, + remove_postproc, + decision_provider, + ) + + def as_json(self, remove_postproc: bool = False) -> JSON_TYPE: + """Serialize the trace as a JSON-style object + + Parameters + ---------- + remove_postproc : bool = False + If postprocessing instructions are removed + + Returns + ---------- + json: JSON_TYPE + The JSON-style object + """ + obj = _ffi_api.TraceAsJSON(self, remove_postproc) # type: ignore # pylint: disable=no-member + return _json_from_tvm(obj) + + def as_python(self, remove_postproc: bool = False) -> List[str]: + """Serialize the trace as a sequence of python statements + + Parameters + ---------- + remove_postproc : bool = False + If postprocessing instructions are removed + + Returns + ---------- + py_stmts: List[str] + A sequence of python statements + """ + return _ffi_api.TraceAsPython(self, remove_postproc) # type: ignore # pylint: disable=no-member + + def with_decision( + self, + inst: Instruction, + decision: DECISION_TYPE, + remove_postproc: bool, + ) -> "Trace": + """Create a new trace with an instruction whose decision is changed, + assuming this instruction exists in the resulting trace + + Parameters + ---------- + inst : Instruction + The instruction whose decision is to be changed + decision : DECISION_TYPE + The decision to be changed to + remove_postproc : bool + If postprocessing instructions are removed + + Returns + ---------- Review comment: Ditto ########## File path: include/tvm/tir/schedule/schedule.h ########## @@ -180,7 +180,8 @@ class ScheduleNode : public runtime::Object { virtual void RemoveRV(const ExprRV& expr_rv) = 0; public: - /******** Block/Loop relation ********/ + /******** Schedule: Sampling ********/ Review comment: Is this (and some following) a placeholder for other schedule primitives? It would be better to put a TODO to make it clearer. ########## File path: python/tvm/tir/schedule/trace.py ########## @@ -0,0 +1,260 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you 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. +"""An execution trace of a scheduling program""" +from typing import TYPE_CHECKING, Any, Callable, Dict, List, Optional + +from tvm._ffi import register_object as _register_object +from tvm.runtime import Object + +from ...ir import Array, Map +from ...runtime import String +from ..expr import FloatImm, IntImm +from . import _ffi_api +from .instruction import ATTR_TYPE, INPUT_RV_TYPE, Instruction + +if TYPE_CHECKING: + from .schedule import ExprRV, Schedule + + +DECISION_TYPE = Any +JSON_TYPE = Any + + +def _json_from_tvm(obj): + if obj is None: + return None + if isinstance(obj, Array): + return [_json_from_tvm(i) for i in obj] + if isinstance(obj, Map): + return {_json_from_tvm(k): _json_from_tvm(v) for k, v in obj.items()} + if isinstance(obj, String): + return str(obj) + if isinstance(obj, (IntImm, FloatImm)): + return obj.value + raise TypeError("Not supported type: " + str(type(obj))) + + +@_register_object("tir.Trace") +class Trace(Object): + """An execution trace of a scheduling program. + + A trace has two parts: + 1) The instructions invoked so far + 2) The random decisions made upon those instructions, if any + + A trace can be serialized to: + 1) Roundtrippable JSON format: can be saved to file and loaded back + 2) Python syntax: allows users to copy-paste the trace to reproduce the scheduling process + + A trace can be applied to a TensorIR schedule by re-applying all its instructions possibly with + their decisions accordingly. Re-sampling is invoked if a sampling instruction doesn't have its + corresponding decision; Otherwise the existing decision will be reused accordingly. + + Attributes + ---------- + insts : List[Instruction] + The instructions invoked so far in the program execution + decisions : Dict[Instruction, DECISION_TYPE] + The random decisions made upon those instructions + """ + + insts: List[Instruction] + decisions: Dict[Instruction, DECISION_TYPE] + + def __init__( + self, + insts: List[Instruction], + decisions: Dict[Instruction, DECISION_TYPE], + ) -> None: + """Constructor + + Parameters + ---------- + insts : List[Instruction] + The instructions invoked so far in the program execution + decisions : Dict[Instruction, DECISION_TYPE] + The random decisions made upon those instructions + """ + self.__init_handle_by_constructor__( + _ffi_api.Trace, # type: ignore # pylint: disable=no-member + insts, + decisions, + ) + + def get_decision(self, inst: Instruction) -> Optional[DECISION_TYPE]: + """Retrieve the decision made on a specific instruction + + Parameters + ---------- + insts : Instruction + The instruction whose decision is to be retrieved + + Returns + ---------- + decision : Optional[DECISION_TYPE] + The corresponding decision; None if there is no decision made on the instruction + """ + return _ffi_api.TraceGetDecision(self, inst) # type: ignore # pylint: disable=no-member + + def append( + self, + inst: Instruction, + decision: Optional[DECISION_TYPE] = None, + ) -> None: + """Append a new instruction to the trace + + Parameters + ---------- + insts : Instruction + The new instruction to be appended + decision : Optional[DECISION_TYPE] = None + The random decision made on this instruction + """ + _ffi_api.TraceAppend(self, inst, decision) # type: ignore # pylint: disable=no-member + + def pop(self) -> Optional[Instruction]: + """Remove the last instruction, along with the decision made on that instruction, if any + + Returns + ---------- + popped_inst : Instruction + Returns the instruction removed; NullOpt if the trace is empty + """ + return _ffi_api.TracePop(self) # type: ignore # pylint: disable=no-member + + def apply_to_schedule( + self, + sch: "Schedule", + remove_postproc: bool, + decision_provider: Optional[ + Callable[ + [Instruction, List[INPUT_RV_TYPE], List[ATTR_TYPE], DECISION_TYPE], DECISION_TYPE + ] + ] = None, + ) -> None: + """Apply the trace to a TensorIR schedule + + Parameters + ---------- + sch : Schedule + The schedule to be applied onto + remove_postproc : bool + If postprocessing instructions are removed + decision_provider: Optional[Callable] = None + A callback that allows users to mutate decisions on the fly when applying instructions. + The signature of the callback is: + - The 1st argument: The instruction + - The 2nd argument: The input random variables + - The 3rd argument: The attributes + - The 4th argument: The decision + - Return: A new decision + """ + _ffi_api.TraceApplyToSchedule( # type: ignore # pylint: disable=no-member + self, + sch, + remove_postproc, + decision_provider, + ) + + def as_json(self, remove_postproc: bool = False) -> JSON_TYPE: + """Serialize the trace as a JSON-style object + + Parameters + ---------- + remove_postproc : bool = False + If postprocessing instructions are removed + + Returns + ---------- Review comment: ```suggestion ------- ``` ########## File path: include/tvm/tir/schedule/instruction.h ########## @@ -0,0 +1,288 @@ +/* + * Licensed to the Apache Software Foundation (ASF) under one + * or more contributor license agreements. See the NOTICE file + * distributed with this work for additional information + * regarding copyright ownership. The ASF licenses this file + * to you 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. + */ +#ifndef TVM_TIR_SCHEDULE_INSTRUCTION_H_ +#define TVM_TIR_SCHEDULE_INSTRUCTION_H_ + +#include <tvm/node/reflection.h> + +#include <utility> + +namespace tvm { + +// Forward declaration +template <typename, typename> +class AttrRegistry; + +namespace tir { + +// Forward declaration +class Schedule; + +/*! + * \brief Type of the functor that applies the instruction to a TensorIR schedule + * \param sch The schedule to be applied on + * \param inputs The input random variables + * \param attrs Instruction attributes + * \param decision Decisions made on the instruction + * \return The functor returns an array of output random variables + */ +using FInstructionApply = runtime::TypedPackedFunc<Array<ObjectRef>( + Schedule sch, const Array<ObjectRef>& inputs, const Array<ObjectRef>& attrs, + const Optional<ObjectRef>& decision)>; + +/*! + * \brief Type of the functor that converts the instruction to a statement in python syntax + * \param inputs Names of the input random variables + * \param attrs Instruction attributes + * \param decisions Decisions made on the instruction + * \param outputs Names of the output random variables + * \return A string representing the python api call + */ +using FInstructionAsPython = runtime::TypedPackedFunc<String( + const Array<ObjectRef>& inputs, const Array<ObjectRef>& attrs, + const Optional<ObjectRef>& decision, const Array<String>& outputs)>; + +/*! + * \brief Type of the functor that serialize its attributes to JSON + * \param attrs The attributes to be serialized + * \return An array, serialized attributes + * \note This functor is nullable + */ +using FInstructionAttrsAsJSON = runtime::TypedPackedFunc<ObjectRef(Array<ObjectRef> attrs)>; + +/*! + * \brief Type of the functor that deserialize its attributes from JSON + * \param json_attrs The attributes to be serialized + * \return An array, deserialized attributes + * \note This functor is nullable + */ +using FInstructionAttrsFromJSON = runtime::TypedPackedFunc<Array<ObjectRef>(ObjectRef json_attrs)>; + +/*! + * \brief Kind of an instruction, e.g. Split, Reorder, etc. + * Besides the name, every kind of instruction has its own properties, including: + * 1) A boolean indicating if the instruction is pure, i.e. change nothing in the schedule state + * 2) A functor that applies the instruction to a TensorIR schedule + * 3) A functor that converts the instruction to a statement in python syntax + * 4) A functor that serialize its attributes to JSON + * 5) A functor that deserialize its attributes from JSON + * + * Unlike `tvm::OpNode`, `InstructionKindNode` doesn't support unstructured properties, + * mainly because there is no such usecase yet to add any other property. + */ +class InstructionKindNode : public runtime::Object { + public: + /*! \brief The name of a kind of instructions */ + String name; + /*! + * \brief Indicates if the instruction is pure, i.e. removing it alone doesn't mutate the schedule + * state. For example, the instruction `GetBlock` is pure because it changes + * nothing, while `ComputeInline` is not because removing it leads to a different resulting + * schedule. + */ + bool is_pure{false}; + /*! \brief A functor that applies the instruction to a TensorIR schedule */ + FInstructionApply f_apply_to_schedule{nullptr}; + /*! \brief A functor that converts the instruction to a statement in python syntax */ + FInstructionAsPython f_as_python{nullptr}; + /*! + * \brief A functor that serialize its attributes to JSON + * \note If the functor is null, it means no conversion is needed + */ + FInstructionAttrsAsJSON f_attrs_as_json{nullptr}; + /*! + * \brief A functor that deserialize its attributes from JSON + * \note If the functor is null, it means no conversion is needed + */ + FInstructionAttrsFromJSON f_attrs_from_json{nullptr}; + + void VisitAttrs(tvm::AttrVisitor* v) { + v->Visit("name", &name); + v->Visit("_is_pure", &is_pure); + // not visited: f_apply_to_schedule + // not visited: f_as_python + // not visited: f_attrs_as_json + // not visited: f_attrs_from_json + } + + static constexpr const char* _type_key = "tir.InstructionKind"; + TVM_DECLARE_FINAL_OBJECT_INFO(InstructionKindNode, runtime::Object); +}; + +/*! + * \brief Managed reference to InstructionKindNode + * \sa InstructionKindNode + */ +class InstructionKind : public runtime::ObjectRef { + public: + /*! + * \brief Retrieve an InstructionKind using its name + * \param name The registered name of the InstructionKind + * \return The InstructionKind retrieved + */ + static InstructionKind Get(const String& name); + TVM_DEFINE_OBJECT_REF_METHODS(InstructionKind, runtime::ObjectRef, InstructionKindNode); +}; + +/*! \brief Schedule instructions each corresponds to a schedule primitive */ +class InstructionNode : public runtime::Object { + public: + /*! \brief The kind of the instruction */ + InstructionKind kind; + /*! + * \brief The input random variables of the instruction, and the type of each element can be one + * of the following: + * - BlockRV + * - LoopRV + * - ExprRV + * - FloatImm + * - IntImm + * - String + * - null pointer + */ + Array<ObjectRef> inputs; + /*! + * \brief The attributes of the instruction. Similar to attributes of an operator, + * attributes of an instruction are arbitrary constant metadata required by the instructions. + * For example, the name of the block to be retrieved in `GetBlock`. + */ + Array<ObjectRef> attrs; + /*! \brief The output random variables of the instruction, and the type of each element can be one + * of the following: + * - BlockRV + * - LoopRV + * - ExprRV, atomic variables only, won't be constants or composite PrimExpr + */ + Array<ObjectRef> outputs; + + void VisitAttrs(tvm::AttrVisitor* v) { + v->Visit("kind", &kind); + v->Visit("inputs", &inputs); + v->Visit("attrs", &attrs); + v->Visit("outputs", &outputs); + } + + static constexpr const char* _type_key = "tir.Instruction"; + TVM_DECLARE_FINAL_OBJECT_INFO(InstructionNode, runtime::Object); +}; + +/*! + * \brief Managed reference to InstructionNode + * \sa InstructionNode + */ +class Instruction : public runtime::ObjectRef { + public: + /*! + * \brief Constructor + * \param kind The kind of the instruction + * \param inputs The input random variables of the instruction + * \param attrs The attributes of the instruction + * \param outputs The output random variables of the instruction + */ + explicit Instruction(InstructionKind kind, Array<ObjectRef> inputs, Array<ObjectRef> attrs, + Array<ObjectRef> outputs); + + TVM_DEFINE_OBJECT_REF_METHODS(Instruction, runtime::ObjectRef, InstructionNode); +}; + +/*! + * \brief A helper macro to register InstructionKind, only used in `TVM_REGISTER_INST_KIND` + * \note This macro is not user-facing. + * \sa TVM_REGISTER_INST_KIND + */ +#define TVM_INST_KIND_REGISTER_VAR_DEF \ + static DMLC_ATTRIBUTE_UNUSED ::tvm::tir::InstructionKindRegEntry& __make_##InstructionKind + +/*! + * \brief Register an InstructionKind + * \param InstructionKindName The name of the InstructionKind + * + * Example: + * + * \code + * + * TVM_REGISTER_INST_KIND("ComputeInline") + * .set_is_pure(false) + * .set_apply_to_schedule(ApplyToSchedule) + * .set_attrs_as_json(AttrsAsJSON) + * .set_attrs_from_json(AttrsFromJSON) + * .set_as_python(AsPython); + * + * \endcode + */ +#define TVM_REGISTER_INST_KIND(InstructionKindName) \ + TVM_STR_CONCAT(TVM_INST_KIND_REGISTER_VAR_DEF, __COUNTER__) = \ + ::tvm::tir::InstructionKindRegEntry::RegisterOrGet(InstructionKindName).set_name() + +/*! \brief An entry in the registry of InstructionKind */ +class InstructionKindRegEntry { + public: + static InstructionKindRegEntry& RegisterOrGet(const String& name); + + InstructionKindRegEntry& set_name() { + get_mutable()->name = this->name; + return *this; + } + + InstructionKindRegEntry& set_is_pure(bool is_pure) { + get_mutable()->is_pure = is_pure; + return *this; + } + + InstructionKindRegEntry& set_apply_to_schedule(FInstructionApply f_apply_to_schedule) { + get_mutable()->f_apply_to_schedule = std::move(f_apply_to_schedule); + return *this; + } + + InstructionKindRegEntry& set_as_python(FInstructionAsPython f_as_python) { + get_mutable()->f_as_python = std::move(f_as_python); + return *this; + } + + InstructionKindRegEntry& set_attrs_as_json(FInstructionAttrsAsJSON f_attrs_as_json) { + get_mutable()->f_attrs_as_json = std::move(f_attrs_as_json); + return *this; + } + + InstructionKindRegEntry& set_attrs_from_json(FInstructionAttrsFromJSON f_attrs_from_json) { + get_mutable()->f_attrs_from_json = std::move(f_attrs_from_json); + return *this; + } + + private: + /*! \brief Private constructor, used only by AttrRegistry */ + explicit InstructionKindRegEntry(uint32_t reg_index); + /*! \brief Get the mutable reference to the internal InstructionKind */ + InstructionKindNode* get_mutable() const { + return const_cast<InstructionKindNode*>(inst_kind_.get()); + } + + /*! \brief The name of the registry entry */ + String name; Review comment: ```suggestion String name_; ``` ########## File path: python/tvm/tir/schedule/trace.py ########## @@ -0,0 +1,260 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you 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. +"""An execution trace of a scheduling program""" +from typing import TYPE_CHECKING, Any, Callable, Dict, List, Optional + +from tvm._ffi import register_object as _register_object +from tvm.runtime import Object + +from ...ir import Array, Map +from ...runtime import String +from ..expr import FloatImm, IntImm +from . import _ffi_api +from .instruction import ATTR_TYPE, INPUT_RV_TYPE, Instruction + +if TYPE_CHECKING: + from .schedule import ExprRV, Schedule + + +DECISION_TYPE = Any +JSON_TYPE = Any + + +def _json_from_tvm(obj): + if obj is None: + return None + if isinstance(obj, Array): + return [_json_from_tvm(i) for i in obj] + if isinstance(obj, Map): + return {_json_from_tvm(k): _json_from_tvm(v) for k, v in obj.items()} + if isinstance(obj, String): + return str(obj) + if isinstance(obj, (IntImm, FloatImm)): + return obj.value + raise TypeError("Not supported type: " + str(type(obj))) + + +@_register_object("tir.Trace") +class Trace(Object): + """An execution trace of a scheduling program. + + A trace has two parts: + 1) The instructions invoked so far + 2) The random decisions made upon those instructions, if any + + A trace can be serialized to: + 1) Roundtrippable JSON format: can be saved to file and loaded back + 2) Python syntax: allows users to copy-paste the trace to reproduce the scheduling process + + A trace can be applied to a TensorIR schedule by re-applying all its instructions possibly with + their decisions accordingly. Re-sampling is invoked if a sampling instruction doesn't have its + corresponding decision; Otherwise the existing decision will be reused accordingly. + + Attributes + ---------- + insts : List[Instruction] + The instructions invoked so far in the program execution + decisions : Dict[Instruction, DECISION_TYPE] + The random decisions made upon those instructions + """ + + insts: List[Instruction] + decisions: Dict[Instruction, DECISION_TYPE] + + def __init__( + self, + insts: List[Instruction], + decisions: Dict[Instruction, DECISION_TYPE], + ) -> None: + """Constructor + + Parameters + ---------- + insts : List[Instruction] + The instructions invoked so far in the program execution + decisions : Dict[Instruction, DECISION_TYPE] + The random decisions made upon those instructions + """ + self.__init_handle_by_constructor__( + _ffi_api.Trace, # type: ignore # pylint: disable=no-member + insts, + decisions, + ) + + def get_decision(self, inst: Instruction) -> Optional[DECISION_TYPE]: + """Retrieve the decision made on a specific instruction + + Parameters + ---------- + insts : Instruction + The instruction whose decision is to be retrieved + + Returns + ---------- + decision : Optional[DECISION_TYPE] + The corresponding decision; None if there is no decision made on the instruction + """ + return _ffi_api.TraceGetDecision(self, inst) # type: ignore # pylint: disable=no-member + + def append( + self, + inst: Instruction, + decision: Optional[DECISION_TYPE] = None, + ) -> None: + """Append a new instruction to the trace + + Parameters + ---------- + insts : Instruction + The new instruction to be appended + decision : Optional[DECISION_TYPE] = None + The random decision made on this instruction + """ + _ffi_api.TraceAppend(self, inst, decision) # type: ignore # pylint: disable=no-member + + def pop(self) -> Optional[Instruction]: + """Remove the last instruction, along with the decision made on that instruction, if any + + Returns + ---------- + popped_inst : Instruction + Returns the instruction removed; NullOpt if the trace is empty + """ + return _ffi_api.TracePop(self) # type: ignore # pylint: disable=no-member + + def apply_to_schedule( + self, + sch: "Schedule", + remove_postproc: bool, + decision_provider: Optional[ + Callable[ + [Instruction, List[INPUT_RV_TYPE], List[ATTR_TYPE], DECISION_TYPE], DECISION_TYPE + ] + ] = None, + ) -> None: + """Apply the trace to a TensorIR schedule + + Parameters + ---------- + sch : Schedule + The schedule to be applied onto + remove_postproc : bool + If postprocessing instructions are removed + decision_provider: Optional[Callable] = None + A callback that allows users to mutate decisions on the fly when applying instructions. + The signature of the callback is: + - The 1st argument: The instruction + - The 2nd argument: The input random variables + - The 3rd argument: The attributes + - The 4th argument: The decision + - Return: A new decision + """ + _ffi_api.TraceApplyToSchedule( # type: ignore # pylint: disable=no-member + self, + sch, + remove_postproc, + decision_provider, + ) + + def as_json(self, remove_postproc: bool = False) -> JSON_TYPE: + """Serialize the trace as a JSON-style object + + Parameters + ---------- + remove_postproc : bool = False + If postprocessing instructions are removed + + Returns + ---------- + json: JSON_TYPE + The JSON-style object + """ + obj = _ffi_api.TraceAsJSON(self, remove_postproc) # type: ignore # pylint: disable=no-member + return _json_from_tvm(obj) + + def as_python(self, remove_postproc: bool = False) -> List[str]: + """Serialize the trace as a sequence of python statements + + Parameters + ---------- + remove_postproc : bool = False + If postprocessing instructions are removed + + Returns + ---------- Review comment: Ditto ########## File path: python/tvm/tir/schedule/trace.py ########## @@ -0,0 +1,260 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you 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. +"""An execution trace of a scheduling program""" +from typing import TYPE_CHECKING, Any, Callable, Dict, List, Optional + +from tvm._ffi import register_object as _register_object +from tvm.runtime import Object + +from ...ir import Array, Map +from ...runtime import String +from ..expr import FloatImm, IntImm +from . import _ffi_api +from .instruction import ATTR_TYPE, INPUT_RV_TYPE, Instruction + +if TYPE_CHECKING: + from .schedule import ExprRV, Schedule + + +DECISION_TYPE = Any +JSON_TYPE = Any + + +def _json_from_tvm(obj): + if obj is None: + return None + if isinstance(obj, Array): + return [_json_from_tvm(i) for i in obj] + if isinstance(obj, Map): + return {_json_from_tvm(k): _json_from_tvm(v) for k, v in obj.items()} + if isinstance(obj, String): + return str(obj) + if isinstance(obj, (IntImm, FloatImm)): + return obj.value + raise TypeError("Not supported type: " + str(type(obj))) + + +@_register_object("tir.Trace") +class Trace(Object): + """An execution trace of a scheduling program. + + A trace has two parts: + 1) The instructions invoked so far + 2) The random decisions made upon those instructions, if any + + A trace can be serialized to: + 1) Roundtrippable JSON format: can be saved to file and loaded back + 2) Python syntax: allows users to copy-paste the trace to reproduce the scheduling process + + A trace can be applied to a TensorIR schedule by re-applying all its instructions possibly with + their decisions accordingly. Re-sampling is invoked if a sampling instruction doesn't have its + corresponding decision; Otherwise the existing decision will be reused accordingly. + + Attributes + ---------- + insts : List[Instruction] + The instructions invoked so far in the program execution + decisions : Dict[Instruction, DECISION_TYPE] + The random decisions made upon those instructions + """ + + insts: List[Instruction] + decisions: Dict[Instruction, DECISION_TYPE] + + def __init__( + self, + insts: List[Instruction], + decisions: Dict[Instruction, DECISION_TYPE], + ) -> None: + """Constructor + + Parameters + ---------- + insts : List[Instruction] + The instructions invoked so far in the program execution + decisions : Dict[Instruction, DECISION_TYPE] + The random decisions made upon those instructions + """ + self.__init_handle_by_constructor__( + _ffi_api.Trace, # type: ignore # pylint: disable=no-member + insts, + decisions, + ) + + def get_decision(self, inst: Instruction) -> Optional[DECISION_TYPE]: + """Retrieve the decision made on a specific instruction + + Parameters + ---------- + insts : Instruction + The instruction whose decision is to be retrieved + + Returns + ---------- + decision : Optional[DECISION_TYPE] + The corresponding decision; None if there is no decision made on the instruction + """ + return _ffi_api.TraceGetDecision(self, inst) # type: ignore # pylint: disable=no-member + + def append( + self, + inst: Instruction, + decision: Optional[DECISION_TYPE] = None, + ) -> None: + """Append a new instruction to the trace + + Parameters + ---------- + insts : Instruction + The new instruction to be appended + decision : Optional[DECISION_TYPE] = None + The random decision made on this instruction + """ + _ffi_api.TraceAppend(self, inst, decision) # type: ignore # pylint: disable=no-member + + def pop(self) -> Optional[Instruction]: + """Remove the last instruction, along with the decision made on that instruction, if any + + Returns + ---------- + popped_inst : Instruction + Returns the instruction removed; NullOpt if the trace is empty + """ + return _ffi_api.TracePop(self) # type: ignore # pylint: disable=no-member + + def apply_to_schedule( + self, + sch: "Schedule", + remove_postproc: bool, + decision_provider: Optional[ + Callable[ + [Instruction, List[INPUT_RV_TYPE], List[ATTR_TYPE], DECISION_TYPE], DECISION_TYPE + ] + ] = None, + ) -> None: + """Apply the trace to a TensorIR schedule + + Parameters + ---------- + sch : Schedule + The schedule to be applied onto + remove_postproc : bool + If postprocessing instructions are removed + decision_provider: Optional[Callable] = None + A callback that allows users to mutate decisions on the fly when applying instructions. + The signature of the callback is: + - The 1st argument: The instruction + - The 2nd argument: The input random variables + - The 3rd argument: The attributes + - The 4th argument: The decision + - Return: A new decision + """ + _ffi_api.TraceApplyToSchedule( # type: ignore # pylint: disable=no-member + self, + sch, + remove_postproc, + decision_provider, + ) + + def as_json(self, remove_postproc: bool = False) -> JSON_TYPE: + """Serialize the trace as a JSON-style object + + Parameters + ---------- + remove_postproc : bool = False + If postprocessing instructions are removed + + Returns + ---------- + json: JSON_TYPE + The JSON-style object + """ + obj = _ffi_api.TraceAsJSON(self, remove_postproc) # type: ignore # pylint: disable=no-member + return _json_from_tvm(obj) + + def as_python(self, remove_postproc: bool = False) -> List[str]: + """Serialize the trace as a sequence of python statements + + Parameters + ---------- + remove_postproc : bool = False + If postprocessing instructions are removed + + Returns + ---------- + py_stmts: List[str] + A sequence of python statements + """ + return _ffi_api.TraceAsPython(self, remove_postproc) # type: ignore # pylint: disable=no-member + + def with_decision( + self, + inst: Instruction, + decision: DECISION_TYPE, + remove_postproc: bool, + ) -> "Trace": + """Create a new trace with an instruction whose decision is changed, + assuming this instruction exists in the resulting trace + + Parameters + ---------- + inst : Instruction + The instruction whose decision is to be changed + decision : DECISION_TYPE + The decision to be changed to + remove_postproc : bool + If postprocessing instructions are removed + + Returns + ---------- + trace: Trace + The new trace with the decision changed + """ + return _ffi_api.TraceWithDecision( # type: ignore # pylint: disable=no-member + self, + inst, + decision, + remove_postproc, + ) + + def simplified(self, remove_postproc: bool) -> "Trace": + """Simplify the trace with dead-code elimination + + Parameters + ---------- + remove_postproc : bool + If postprocessing instructions are removed + + Returns + ---------- Review comment: Ditto ########## File path: tests/python/unittest/test_tir_schedule_trace.py ########## @@ -0,0 +1,248 @@ +# Licensed to the Apache Software Foundation (ASF) under one +# or more contributor license agreements. See the NOTICE file +# distributed with this work for additional information +# regarding copyright ownership. The ASF licenses this file +# to you 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. +# pylint: disable=missing-function-docstring,missing-module-docstring +# mypy: ignore-errors +import tvm +from tvm.script import ty +from tvm import tir +from tvm.tir.schedule import BlockRV, Instruction, InstructionKind, LoopRV, Trace + + +# pylint: disable=no-member,invalid-name,unused-variable + + [email protected] +def elementwise(a: ty.handle, c: ty.handle) -> None: + A = tir.match_buffer(a, (128, 128)) + B = tir.alloc_buffer((128, 128)) + C = tir.match_buffer(c, (128, 128)) + with tir.block([128, 128], "B") as [vi, vj]: + B[vi, vj] = A[vi, vj] * 2.0 + with tir.block([128, 128], "C") as [vi, vj]: + C[vi, vj] = B[vi, vj] + 1.0 + + [email protected] +def elementwise_inlined(a: ty.handle, c: ty.handle) -> None: + A = tir.match_buffer(a, (128, 128)) + C = tir.match_buffer(c, (128, 128)) + with tir.block([128, 128], "C") as [vi, vj]: + C[vi, vj] = A[vi, vj] * 2.0 + 1.0 + + +# pylint: enable=no-member,invalid-name,unused-variable + + +def _make_get_block(name, output): + return Instruction( + kind=InstructionKind.get("GetBlock"), + inputs=[], + attrs=[name, "main"], + outputs=[output], + ) + + +def _make_get_loops(input, outputs): # pylint: disable=redefined-builtin + return Instruction( + kind=InstructionKind.get("GetLoops"), + inputs=[input], + attrs=[], + outputs=outputs, + ) + + +def _make_compute_inline(input): # pylint: disable=redefined-builtin + return Instruction( + kind=InstructionKind.get("ComputeInline"), + inputs=[input], + attrs=[], + outputs=[], + ) + + +def _make_enter_postproc(): + return Instruction( + kind=InstructionKind.get("EnterPostproc"), + inputs=[], + attrs=[], + outputs=[], + ) + + +def _make_trace_1(b0, l1, l2): # pylint: disable=invalid-name + return Trace( + insts=[ + _make_get_block(name="block", output=b0), + _make_get_loops(input=b0, outputs=[l1, l2]), + ], + decisions={}, + ) + + +def _make_trace_2(b0): # pylint: disable=invalid-name + return Trace( + insts=[ + _make_get_block(name="B", output=b0), + _make_compute_inline(input=b0), + ], + decisions={}, + ) + + +def _make_trace_3(b0, b1, add_postproc): # pylint: disable=invalid-name + if add_postproc: + insts = [ + _make_get_block(name="B", output=b0), + _make_compute_inline(input=b0), + _make_get_block(name="C", output=b1), + _make_enter_postproc(), + _make_compute_inline(input=b1), + ] + else: + insts = [ + _make_get_block(name="B", output=b0), + _make_compute_inline(input=b0), + _make_get_block(name="C", output=b1), + ] + return Trace(insts=insts, decisions={}) + + +def test_trace_construct_1(): + trace = _make_trace_1(BlockRV(), LoopRV(), LoopRV()) + assert str(trace) == "\n".join( + ( + 'b0 = sch.get_block(name="block", func_name="main")', + "l1, l2 = sch.get_loops(block=b0)", + ) + ) + assert len(trace.insts) == 2 + assert len(trace.decisions) == 0 + + +def test_trace_construct_get_decision_1(): + trace = _make_trace_1(BlockRV(), LoopRV(), LoopRV()) + assert trace.get_decision(trace.insts[0]) is None + assert trace.get_decision(trace.insts[1]) is None + + +def test_trace_construct_append_1(): + trace = _make_trace_1(BlockRV(), LoopRV(), LoopRV()) + trace.append(inst=_make_get_block("block2", BlockRV())) + assert str(trace) == "\n".join( + ( + 'b0 = sch.get_block(name="block", func_name="main")', + "l1, l2 = sch.get_loops(block=b0)", + 'b3 = sch.get_block(name="block2", func_name="main")', + ) + ) + + +def test_trace_construct_pop_1(): + trace = _make_trace_1(BlockRV(), LoopRV(), LoopRV()) + last_inst = trace.insts[-1] + assert trace.pop().same_as(last_inst) + assert str(trace) == 'b0 = sch.get_block(name="block", func_name="main")' + + +def test_trace_construct_pop_2(): + trace = Trace([], {}) + assert str(trace) == "" + assert trace.pop() is None + assert str(trace) == "" + + +def test_trace_apply_to_schedule(): + trace = _make_trace_2(BlockRV()) + sch = tir.Schedule(elementwise, debug_mode=True) + trace.apply_to_schedule(sch, remove_postproc=False, decision_provider=None) + tvm.ir.assert_structural_equal(elementwise_inlined, sch.mod["main"]) + + +def test_trace_as_json_1(): + trace = _make_trace_1(BlockRV(), LoopRV(), LoopRV()) + obj = trace.as_json() + assert obj == [ + [ + ["GetBlock", [], ["block", "main"], ["b0"]], + ["GetLoops", ["b0"], [], ["l1", "l2"]], + ], + [], + ] + + +def test_trace_simplified_1(): + trace = _make_trace_3(BlockRV(), BlockRV(), add_postproc=True) + assert str(trace) == "\n".join( + ( + 'b0 = sch.get_block(name="B", func_name="main")', + "sch.compute_inline(block=b0)", + 'b1 = sch.get_block(name="C", func_name="main")', + "sch.enter_postproc()", + "sch.compute_inline(block=b1)", + ) + ) + trace = trace.simplified(remove_postproc=True) + assert str(trace) == "\n".join( + ( + 'b0 = sch.get_block(name="B", func_name="main")', + "sch.compute_inline(block=b0)", + ) + ) + + +def test_trace_simplified_2(): + trace = _make_trace_3(BlockRV(), BlockRV(), add_postproc=True) + assert str(trace) == "\n".join( + ( + 'b0 = sch.get_block(name="B", func_name="main")', + "sch.compute_inline(block=b0)", + 'b1 = sch.get_block(name="C", func_name="main")', + "sch.enter_postproc()", + "sch.compute_inline(block=b1)", + ) + ) + trace = trace.simplified(remove_postproc=False) + assert str(trace) == "\n".join( + ( + 'b0 = sch.get_block(name="B", func_name="main")', + "sch.compute_inline(block=b0)", + 'b1 = sch.get_block(name="C", func_name="main")', + "sch.enter_postproc()", + "sch.compute_inline(block=b1)", + ) + ) + + +def test_apply_json_to_schedule_1(): + trace = _make_trace_2(BlockRV()) + json_obj = trace.as_json() + sch = tir.Schedule(elementwise, debug_mode=True) + Trace.apply_json_to_schedule(json_obj, sch) + tvm.ir.assert_structural_equal(elementwise_inlined, sch.mod["main"]) + + +if __name__ == "__main__": Review comment: Better to use pytest method to collect tests. -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected]
