# 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=invalid-name, no-member"""VM build logics"""fromtypingimportDict,List,Optional,Unionimporttvmfromtvmimportrelaxfromtvm.ir.moduleimportIRModulefromtvm.tir.functionimportPrimFuncfromtvm.runtimeimportExecutablefrom.import_ffi_api
[文档]classVMExecutable(Executable):"""The virtual machine executable object emitted by the VM compiler or the ExecBuilder."""def__init__(self,mod:tvm.runtime.Module):super().__init__(mod)self._stats=self.mod["stats"]self._as_text=self.mod["as_text"]self._as_python=self.mod["as_python"]
[文档]defstats(self)->str:"""print the detailed statistics of the executable."""returnself._stats()
[文档]defas_text(self)->str:"""print the instructions as text format."""returnself._as_text()
[文档]defas_python(self)->str:"""print the instructions as python program."""returnself._as_python()
def_vmcodegen(builder:"relax.ExecBuilder",mod:tvm.IRModule,exec_mode:str="bytecode",)->tvm.IRModule:"""Running VM codegen. Parameters ---------- builder: relax.ExecBuilder ExecBuilder to collect the vm executable. mod: IRModule The input IRModule to be built. exec_mode: {"bytecode", "compiled"} The execution mode. Return ------ leftover: IRModule Left over IRModule that may contain extra functions. """ifexec_mode=="bytecode":return_ffi_api.VMCodeGen(builder,mod)# type:ignoreifexec_mode=="compiled":return_ffi_api.VMTIRCodeGen(builder,mod)# type: ignoreraiseValueError(f"Unknown exec_mode {exec_mode}")def_auto_attach_system_lib_prefix(tir_mod:tvm.IRModule,target:Optional[tvm.target.Target]=None,system_lib:Optional[bool]=None,):"""Automatically detect system lib req and attach prefix attr"""iftargetisnotNone:host=targetiftarget.hostisNoneelsetarget.hostifsystem_libisNone:system_lib=Falseif"wasm"inhost.attrs.get("mtriple",""):system_lib=Trueifsystem_lib:iftir_mod.get_attr("system_lib_prefix")isNone:returntir_mod.with_attr("system_lib_prefix","")returntir_moddef_vmlink(builder:"relax.ExecBuilder",target:Optional[Union[str,tvm.target.Target]],tir_mod:Optional[tvm.IRModule]=None,tir_pipeline:Optional[Union[str,tvm.transform.Pass]]="default",ext_libs:List[tvm.runtime.Module]=None,params:Optional[Dict[str,list]]=None,*,system_lib:Optional[bool]=None,):""" Internal codegen function to make executable. This function is only used for unit-testing purpoes. Use build instead. Parameters ---------- builder: relax.ExecBuilder Builder used to collect executables. target : Optional[Union[str, tvm.target.Target]] A build target which can have optional host side compilation target. If the target is not specified, the target in the vdevice list will be used. For multi-target compilation, the vdevice should be annotated. tir_mod: IRModule The input TIR IRModule to be linked together. ext_libs: List[tvm.runtime.Module] List of compiled external modules. params: Optional[Dict[str, list]] Extra parameter mappings. Returns ------- ex: tvm.relax.Executable An executable that can be loaded by virtual machine. """ifisinstance(target,str):target=tvm.target.Target(target)ifparamsisNone:params={}ifext_libsisNone:ext_libs=[]lib=Nonerelax_ext_libs=[]tir_ext_libs=[]iftir_modisnotNoneandlen(tir_mod.get_global_vars())>0:tir_mod=_auto_attach_system_lib_prefix(tir_mod,target,system_lib)lib=tvm.tir.build(tir_mod,target=target,pipeline=tir_pipeline)forext_modinext_libs:ifext_mod.is_device_module:tir_ext_libs.append(ext_mod)else:relax_ext_libs.append(ext_mod)iflibisnotNone:formodintir_ext_libs:lib.import_module(mod)eliflen(tir_ext_libs)>0:print("Warning: No TIR module is found, but external modules for TIR are provided.")lib=_ffi_api.VMLink(builder,target,lib,relax_ext_libs,params)# type: ignorereturnVMExecutable(lib)
[文档]defbuild(mod:tvm.IRModule,target:Optional[Union[str,tvm.target.Target]]=None,params:Optional[Dict[str,list]]=None,relax_pipeline:Union[None,str,tvm.transform.Pass]="default",tir_pipeline:Union[None,str,tvm.transform.Pass]="default",exec_mode:str="bytecode",*,system_lib:Optional[bool]=None,)->Executable:""" Build an IRModule to VM executable. Parameters ---------- mod: IRModule The input IRModule to be built. target : Optional[Union[str, tvm.target.Target]] A build target which can have optional host side compilation target. When TVM compiles device specific program such as CUDA, we also need host(CPU) side code to interact with the driver to setup the dimensions and parameters correctly. host is used to specify the host side codegen target. By default, llvm is used if it is enabled, otherwise a stackvm interpreter is used. params: Optional[Dict[str, list]] Parameters for the input IRModule that will be bound. relax_pipeline : str = "default" The Relax compilation pipeline to use. tir_pipelinie : str = "default" The TIR compilation pipeline to use. exec_mode: {"bytecode", "compiled"} The execution mode. system_lib: Optional[bool] Whether to build system lib that is being packed statically and auto registers generated functions to the system. By default auto detects based on the target. Returns ------- ex: tvm.relax.Executable An executable that can be loaded by virtual machine. Example ------- .. code-block:: python class InputModule: @R.function def foo(x: Tensor((3, 4), "float32"), y: Tensor((3, 4), "float32")): z = R.add(x, y) return z mod = InputModule target = tvm.target.Target("llvm", host="llvm") ex = tvm.compile(mod, target) """def_extract_attrs(mod:tvm.IRModule):attrs=dict(mod.attrs)ifmod.attrselse{}ext_libs=attrs.get("external_mods",[])constants=attrs.get("const_name_to_constant",{})returnext_libs,constantsifisinstance(target,str):target=tvm.target.Target(target)ifnotparams:params={}ifrelax_pipelineisnotNone:ifisinstance(relax_pipeline,str):relax_pipeline=relax.get_pipeline(relax_pipeline)iftargetisNone:mod=relax_pipeline(mod)else:withtarget:mod=relax_pipeline(mod)ext_libs,constants=_extract_attrs(mod)params.update(dict(constants))builder=relax.ExecBuilder()mod=_vmcodegen(builder,mod,exec_mode)return_vmlink(builder=builder,target=target,tir_mod=_filter_tir(mod),tir_pipeline=tir_pipeline,ext_libs=ext_libs,params=params,system_lib=system_lib,)