模型编译深度学习推理引擎【免费下载链接】tvmOpen Machine Learning Compiler Framework项目地址https://gitcode.com/gh_mirrors/tv/tvm点击查看免费下载S-TIRSchedulable TensorIR是 Apache TVM 中以可调度 block 为中心的中间表示层而tvm.s_tir.transform模块则集中封装了 S-TIR 专用的全部变换 Pass负责把带 block 结构、可继续调度的 TIR 一步步精简化、规范化并最终降低为不可再调度的普通 TIR。本文以仓库中的 API 参考文档 docs/reference/api/python/s_tir/transform.rst 为骨架逐一梳理该模块暴露的每个变换 Pass 的用途、参数与底层实现并结合默认编译管线展示它们在实际后端生成中的编排顺序帮助读者掌握 S-TIR 到目标代码之间的“后半程”到底发生了什么。模块定位S-TIR 变换在 TVM 编译栈中的位置S-TIR 是 TVM 中“可调度的 TensorIR”其核心特征是语句树中保留着SBlock、SBlockRealize等可调度节点对应 include/tvm/s_tir/stmt.h 与 Python 侧暴露的 python/tvm/s_tir/stmt.py。调度器python/tvm/s_tir/schedule/init.py完成 loop split、fuse、binding 等操作后产生的仍是带 block 结构的 S-TIR此时需要一系列变换 Pass 把它“降级”成普通 TIR再交给后续的向量化、存储重写、寄存器分配等后端处理。tvm.s_tir.transform正是这组降级与优化 Pass 的 Python 命名空间。它的文档入口是docs/reference/api/python/s_tir/transform.rst该文件通过 Sphinx 的automodule指令自动展开模块内所有公开成员tvm.s_tir.transform ------------------- .. automodule:: tvm.s_tir.transform :members: :imported-members: :no-index:因此该文档展开后的实质内容就是python/tvm/s_tir/transform/transform.py中定义的 40 余个工厂函数——每个函数返回一个tvm.transform.Pass对象。这些 Pass 的 C 实现全部位于 src/s_tir/transform如canonicalize_loop.cc、lower_opaque_block.cc、inject_software_pipeline.cc等Python 层通过 FFI 绑定接入python/tvm/s_tir/transform/transform.py全部 Pass 工厂函数定义python/tvm/s_tir/transform/_ffi_api.pytvm_ffi.init_ffi_api(s_tir.transform, __name__)完成与 C 注册表s_tir.transform.*的绑定python/tvm/s_tir/transform/init.py通配导入上述定义并额外从tvm.tirx.transform引入HoistedConditionals、HoistedLetBindings两个位标志类型。使用方式也很统一pass_ tvm.s_tir.transform.CanonicalizeLoop()得到 Pass 对象再放入tvm.ir.transform.Sequential([...])或PassContext中执行。下面按功能族逐组介绍。结构化精简与规范化 Pass这一类 Pass 负责把 S-TIR 中的可调度结构block、block realize、block var逐步消除或规范化是整个降级过程的前置步骤。CanonicalizeLoopdef CanonicalizeLoop(): Canonicalize the loop to start from zero and use trivial step return _ffi_api.CanonicalizeLoop()将循环规范化为从 0 起始、步长为 1trivial step的标准形式为后续 loop partition、向量化等 Pass 提供统一的循环结构。C 实现见 src/s_tir/transform/canonicalize_loop.cc。StmtSimplifydef StmtSimplify(): Simplify schedulable TIR with block constraints and tirx.StmtSimplify options.在 block 约束如 block var 的取值范围存在的情况下简化语句。注意它与tvm.tirx.transform.StmtSimplify不同——前者面向“仍可调度”的 S-TIR会利用 block 的迭代域约束来简化后者面向普通 TIR。实现位于 src/s_tir/transform/stmt_simplify.cc 及配套的 stmt_simplify.h。ConvertSSAdef ConvertSSA(): De-duplicate definitions, including schedulable block iterators, across PrimFuncs.在多个 PrimFunc 之间去重变量与 buffer 定义含可调度的 block 迭代器保证每个定义只出现一次形成严格的 SSA 形式。在 python/tvm/s_tir/init.py 中还有一个相关的顶层函数s_tir.renew_defs(func)其底层对应RenewDefs作用类似 DeepCopy用一组全新的 Var/Buffer 复制出一个行为等价的函数。RenormalizeSplitPatterndef RenormalizeSplitPattern(): Renormalize the split pattern from floordiv(floormod()) to floormod(floordiv())调度中 split 产生的索引计算通常呈现floordiv(floormod(...))的嵌套形式该 Pass 将其重写为floormod(floordiv(...))的规范形态利于后端识别访问模式。实现见 src/s_tir/transform/renormalize_split_pattern.cc。Block 与 Buffer 相关 Pass这一组直接操作 S-TIR 的 block 结构定位分配位置、把 block 变 opaque、按实际访问域收缩 buffer、最终移除 block。PlanAndUpdateBufferAllocationLocationdef PlanAndUpdateBufferAllocationLocation(): Locate the buffer allocation to the exact position (usually is the lca of buffer access). This pass will inject opaque block with alloc_buffers at the allocation site.把 buffer 的分配位置移动到其所有访问点的最近公共祖先LCA处并在该位置注入一个带alloc_buffers的 opaque block。这是把“逻辑上在函数头分配”的 buffer 改为“贴近使用现场分配”的关键一步实现见 src/s_tir/transform/plan_update_buffer_allocation_location.cc。ConvertBlocksToOpaquedef ConvertBlocksToOpaque(): Substitute all the block vars with the PrimExprs they are bound to, indicated by the corresponding iter_values in BlockRealize, and then convert the blocks into opaque ones by removing all the iter_values in BlockRealize and iter_vars in Block.先用BlockRealize中的iter_values替换掉所有 block var然后删除iter_values与iter_vars把每个 block 变成 opaque block。这是 S-TIR 失去“可调度性”的分水岭——opaque block 内部的迭代关系已经被展开调度器不再能对其做 split/fuse 类操作。实现见 src/s_tir/transform/convert_blocks_to_opaque.cc。CompactBufferAllocationdef CompactBufferAllocation(is_strict: bool True): Compact the buffer access region by removing the buffer regions that are not accessed, i.e. narrowing the buffer shape and adjust the access region if necessary. Parameters ---------- is_strict : bool Ensure the compacted shape to be always smaller than the original shape. Otherwise it allows to grow the shape to match actual accessed buffer regions. return _ffi_api.CompactBufferAllocation(is_strict)按“实际被访问的区域”收缩 buffer 形状删除从未被访问的 region必要时同步调整访问索引。参数is_strict控制策略True默认收缩后的形状必须严格小于等于原形状绝不扩容False允许形状增长以完全匹配实际访问域例如某些边界访问需要 padding 时。实现见 src/s_tir/transform/compact_buffer_region.cc。LowerMatchBuffer 与 LowerOpaqueBlockdef LowerMatchBuffer(): Remove match buffers inside the block. Also, it will validate the binding. def LowerOpaqueBlock(): Remove the block to ensure that the TIR can not be scheduled again.LowerMatchBuffer移除 block 内部的 match buffer同时校验 binding 合法性实现见 src/s_tir/transform/lower_match_buffer.ccLowerOpaqueBlock直接删除 opaque block把 block 内的语句展开到外层作用域。执行完这个 Pass 后TIR 中不再有任何 block也就“不可能再被调度”彻底完成了可调度性到普通 TIR 的过渡。实现见 src/s_tir/transform/lower_opaque_block.cc。线程绑定、并行与流水线 Pass这一组面向多线程/多核后端尤其是 GPU处理线程绑定、跨线程归约、共享内存同步、软件流水线、双缓冲等。LiftThreadBinding 与 UnifyThreadBindingdef LiftThreadBinding(): Lift the same thread bindings to their LCA loops. def UnifyThreadBinding(): Unify all the thread bindings for blockIdx.x/y/z, threadIdx.x/y/z, and vthread.x/y/z.LiftThreadBinding把相同的线程绑定提升到它们共同的最外层循环LCA减少重复绑定语句实现见 src/s_tir/transform/lift_thread_binding.ccUnifyThreadBinding统一blockIdx.x/y/z、threadIdx.x/y/z、vthread.x/y/z的绑定实现见 src/s_tir/transform/unify_thread_binding.cc。LowerCrossThreadReduction 与 LowerThreadAllreducedef LowerCrossThreadReduction(): Lower cross-thread reduction from thread bindings to intrinsic function calls. def LowerThreadAllreduce(): Lower cross thread allreduce.LowerCrossThreadReduction把跨线程归约基于线程绑定形式降低为 intrinsic 函数调用见 src/s_tir/transform/lower_cross_thread_reduction.ccLowerThreadAllreduce降低跨线程 allreduce见 src/s_tir/transform/lower_thread_allreduce.cc。两者都服务于 GPU 上多线程协同归约的场景最终配合ThreadSync插入的同步保证正确性。ThreadSyncdef ThreadSync(storage_scope): Insert sync between parallel read/write of shared buffers. Parameters ---------- storage_scope: str The target storage scope. return _ffi_api.ThreadSync(storage_scope)在共享 buffer 的并行读写之间插入同步原语。storage_scope指定作用域仓库中默认管线分别以shared、shared.dyn、warp调用三次见下文管线编排。实现见 src/s_tir/transform/thread_storage_sync.cc。InjectVirtualThread 与 InjectDoubleBufferdef InjectVirtualThread(): Inject virtual thread loops. _ffi.register_object(s_tir.transform.InjectDoubleBufferConfig) class InjectDoubleBufferConfig(_ffi.Object): Config for inject double buffer pass def InjectDoubleBuffer(): Inject double buffer statements.InjectVirtualThread注入虚拟线程循环实现见 src/s_tir/transform/inject_virtual_thread.ccInjectDoubleBuffer注入双缓冲语句并配套注册了InjectDoubleBufferConfig配置对象实现见 src/s_tir/transform/inject_double_buffer.cc。InjectSoftwarePipeline 与 ManifestSharedMemoryLocalStagedef InjectSoftwarePipeline(): Transform annotated loops into pipelined one that parallelize producers and consumers def ManifestSharedMemoryLocalStage(): Add the explicit local stage for the shared memory access on GPU.InjectSoftwarePipeline将带流水线注解的循环转换为生产者/消费者并行的软件流水线实现见 src/s_tir/transform/inject_software_pipeline.ccManifestSharedMemoryLocalStage为 GPU 共享内存访问显式添加 local stage实现见 src/s_tir/transform/manifest_shared_memory_local_stage.cc。LowerInitBlock、DefaultGPUSchedule 与 AnnotateIrregularLoopdef LowerInitBlock(): Lower block init stmt into IfThenElse statements. def DefaultGPUSchedule(): Set default thread bindings for GPU PrimFuncs. def AnnotateIrregularLoop(): Annotate irregular loop mark.LowerInitBlock把 block 的 init 语句降低为IfThenElse实现见 src/s_tir/transform/lower_init_block.ccDefaultGPUSchedule为 GPU PrimFunc 设置默认线程绑定实现见 src/s_tir/transform/default_gpu_schedule.ccAnnotateIrregularLoop为不规则循环如边界不齐的循环打上注解标记便于后端做特殊处理实现见 src/s_tir/transform/annotate_irregular_loop.cc。GPU 特化TensorCore、Async Copy 与内存合并InferFragment、TransformMmaBufferLayout 与 InjectPermutedLayoutdef InferFragment(): Infer the TensorCore fragment information using tensor intrinsics. def TransformMmaBufferLayout(): Transform mma buffer layout def InjectPermutedLayout(): Inject permuted layout in mma这三个 Pass 共同服务于 TensorCore / MMA 指令InferFragment利用 tensor intrinsics 推断 TensorCore 的 fragment 信息A/B/C 矩阵在各线程上的分布见 src/s_tir/transform/tensorcore_infer_fragment.ccTransformMmaBufferLayout变换 MMA 相关的 buffer 布局见 src/s_tir/transform/transform_mma_buffer_layout.ccInjectPermutedLayout在 MMA 中注入置换后的布局见 src/s_tir/transform/inject_permuted_layout.cc。InjectPTXAsyncCopy 与 LowerAsyncDMAdef LowerAsyncDMA(): Lower async DMA to DMA. def InjectPTXAsyncCopy(): Rewrite global to shared memory copy on CUDA with asynchronous copy.LowerAsyncDMA把异步 DMA 原语降低为普通 DMA见 src/s_tir/transform/lower_async_dma.ccInjectPTXAsyncCopy把 CUDA 上 global 到 shared 的拷贝改写为异步拷贝cp.async类指令见 src/s_tir/transform/inject_ptx_async_copy.cc。InjectPTXLDG32def InjectPTXLDG32(enable_inject_ptx_intrinTrue): Inject ptx.ldg.32 intrinsics. Parameters ---------- enable_inject_ptx_intrin : bool If True, inject ptx.ldg.32 intrinsics. return _ffi_api.InjectPTXLDG32(enable_inject_ptx_intrin)注入ptx.ldg.32intrinsic通过 read-only cache 的 32 位加载。参数enable_inject_ptx_intrin控制是否注入默认True。实现见 src/s_tir/transform/inject_ptx_ldg32.cc。仓库默认管线中它会出现两次一次带True在tirx.s_tir.ldg32配置开启时一次不带参数。MergeSharedMemoryAllocationsdef MergeSharedMemoryAllocations(): This pass merges multiple TIR-level shared memory allocations into one allocation.把多个 TIR 层面的共享内存分配合并为一次分配减少__shared__声明的数量与对齐开销实现见 src/s_tir/transform/merge_shared_memory_allocations.cc。表达式提升与代码质量优化 PassHoistIfThenElse 与 HoistExpressiondef HoistIfThenElse(variantNone): Hoist loop-invariant IfThenElse nodes to outside the eligible loops. Parameters ---------- variant : Optional[String] The variant of the pass. variant can have any one of following values [basic, None(Default)]. if variant basic: return _ffi_api.HoistIfThenElseBasic() elif variant is None: return _ffi_api.HoistIfThenElse() else: raise ValueError(wrong variant of HoistIfThenElse, variant) def HoistExpression(): Hoist loop-invariant expressions to outside the eligible loops.HoistIfThenElse把循环不变的条件分支提升到循环外。variant参数支持basic基础变体与默认变体两种实现传入其他值会抛出ValueErrorHoistExpression提升循环不变表达式实现位于 src/s_tir/transform/hoist_expression.cc。该 Pass 内部定义了两个位标志类型HoistedConditionals与HoistedLetBindings如kIfElseStmt、kIfElseExpr、kBooleanExpression、kRequiredByCondition、kBind、kLetExpr它们作为配置位被注册并由 python/tvm/s_tir/transform/init.py 从tvm.tirx.transform再导出供用户控制具体提升哪些结构。LoopPartition_ffi.register_object(s_tir.transform.LoopPartitionConfig) class LoopPartitionConfig(_ffi.Object): Config for loop partition pass def LoopPartition(): Partition loops in the stmt.按边界条件将循环切分为多个子循环如主循环 边界处理循环并注册了LoopPartitionConfig配置对象。实现见 src/s_tir/transform/loop_partition.cc。RewriteUnsafeSelect、RemoveStoreUndef 与 UseAssumeToReduceBranchesdef RewriteUnsafeSelect(): Detect and rewrite unsafe select that contains memory access. def RemoveStoreUndef(): Remove stores of undefined values from the Stmt. def UseAssumeToReduceBranches(): Eliminate layout specific pad branch by overcomputing values for padded region.RewriteUnsafeSelect检测并重写包含内存访问的不安全 select避免条件不成立时仍触发访存见 src/s_tir/transform/rewrite_unsafe_select.ccRemoveStoreUndef删除对未定义值的 store见 src/s_tir/transform/remove_store_undef.ccUseAssumeToReduceBranches通过为 padding 区域“过度计算”值来消除 layout 特有的 padding 分支配合 assume 使用见 src/s_tir/transform/using_assume_to_reduce_branches.cc。DecorateDeviceScopedef DecorateDeviceScope(): Decorate all the functions body as device function.把函数体整体标记为 device function见 src/s_tir/transform/decorate_device_scope.cc。目标感知与诊断 PassVerifyVTCMLimit 与 LowerVtcmAllocdef VerifyVTCMLimit(default_targetNone): Verify if the size of the allocated vtcm memory satisfies the limit. The limit is determined from the vtcm-capacity attribute of the target. Parameters ---------- default_target : Optional[tvm.target.Target] The default target to use if a PrimFunc does not have a target attribute. return _ffi_api.VerifyVTCMLimit(default_target) def LowerVtcmAlloc(): Lower vtcm allocation.VerifyVTCMLimit校验 VTCM片上紧耦合内存分配量是否满足 target 的vtcm-capacity属性限制default_target用于没有 target 属性的 PrimFunc 兜底LowerVtcmAlloc把 VTCM 分配降低为具体目标上的实现见 src/s_tir/transform/lower_vtcm_alloc.cc。这两个 Pass 的顺序有硬性要求VerifyVTCMLimit必须先于LowerVtcmAlloc执行默认管线中有注释明确这一点。InstrumentBoundCheckers 与 InstrumentProfileIntrinsicsdef InstrumentBoundCheckers(): Instruments bound checkers. def InstrumentProfileIntrinsics(): Insert intrinsic calls to instrument function and loop level profiling.InstrumentBoundCheckers插入边界检查器实现见 src/s_tir/transform/bound_checker.ccInstrumentProfileIntrinsics插入 intrinsic 调用以对函数与循环级别做性能剖析见 src/s_tir/transform/profile_instrumentation.cc。RemoveWeightLayoutRewriteBlockdef RemoveWeightLayoutRewriteBlock(skip_tensor_rewriteFalse): Remove weight layout rewrite block before benchmarking during tuning stage. Parameters ---------- skip_tensor_rewrite : bool If True, exact rewrite of Tensor, according to the given index map, will be skipped. return _ffi_api.RemoveWeightLayoutRewriteBlock(skip_tensor_rewrite)在调优tuning阶段的 benchmark 之前移除权重 layout 重写 blockskip_tensor_rewriteTrue时跳过根据 index map 对 Tensor 的精确重写。实现见 src/s_tir/transform/remove_weight_layout_rewrite_block.cc。实战这些 Pass 在默认 S-TIR 管线中的编排仅了解单个 Pass 不够真正的价值在于它们的组合。仓库中 python/tvm/s_tir/pipeline.py 的default_s_tir_pipeline()以tvm.ir.transform.Sequential的形式编排了上述大部分 Pass展示了 S-TIR 后端完整的降级顺序节选核心部分passes [ s_tir.transform.CanonicalizeLoop(), s_tir.transform.LowerCrossThreadReduction(), s_tir.transform.LowerInitBlock(), s_tir.transform.PlanAndUpdateBufferAllocationLocation(), s_tir.transform.ConvertBlocksToOpaque(), s_tir.transform.LiftThreadBinding(), s_tir.transform.ManifestSharedMemoryLocalStage(), s_tir.transform.CompactBufferAllocation(), s_tir.transform.LowerAutoCopy(), s_tir.transform.UnifyThreadBinding(), s_tir.transform.LowerMatchBuffer(), s_tir.transform.StmtSimplify(), s_tir.transform.InjectPermutedLayout(), s_tir.transform.AnnotateIrregularLoop(), s_tir.transform.InjectSoftwarePipeline(), s_tir.transform.TransformMmaBufferLayout(), s_tir.transform.LowerOpaqueBlock(), tirx.transform.FlattenBuffer(), tirx.transform.NarrowDataType(32), s_tir.transform.LoopPartition(), s_tir.transform.InjectVirtualThread(), s_tir.transform.InjectDoubleBuffer(), ... s_tir.transform.HoistIfThenElse(), s_tir.transform.RenormalizeSplitPattern(), s_tir.transform.RewriteUnsafeSelect(), ... s_tir.transform.VerifyVTCMLimit(), s_tir.transform.LowerVtcmAlloc(), ... s_tir.transform.ThreadSync(shared), s_tir.transform.ThreadSync(shared.dyn), s_tir.transform.ThreadSync(warp), s_tir.transform.InferFragment(), s_tir.transform.LowerThreadAllreduce(), ... s_tir.transform.MergeSharedMemoryAllocations(), tirx.transform.SplitHostDevice(), ... ]从这份编排可以清晰读出一套完整的降级叙事规范与精化CanonicalizeLoop→LowerInitBlock→StmtSimplify先让结构规整分配定位与 block 降级PlanAndUpdateBufferAllocationLocation→ConvertBlocksToOpaque→CompactBufferAllocation→LowerMatchBuffer→LowerOpaqueBlock逐步把可调度 block 变成普通语句GPU 特化LiftThreadBinding/UnifyThreadBinding、ManifestSharedMemoryLocalStage、InjectSoftwarePipeline、InjectPermutedLayout/TransformMmaBufferLayout、ThreadSync/InferFragment/LowerThreadAllreduce、MergeSharedMemoryAllocations后端收尾VerifyVTCMLimit必须先于LowerVtcmAlloc随后SplitHostDevice把 host/device 代码分开。此外管线中还有若干受配置开关控制的 Pass展示了这些 Pass 与PassContext配置的联动方式if not bool(config.get(tirx.disable_storage_rewrite, False)): passes.append(tirx.transform.StorageRewrite()) if config.get(tirx.use_async_copy, False): passes.append(s_tir.transform.LowerAsyncDMA()) ... if bool(config.get(tirx.instrument_bound_checkers, False)): passes.append(s_tir.transform.InstrumentBoundCheckers()) if bool(config.get(tirx.s_tir.ldg32, False)): passes.append(s_tir.transform.InjectPTXLDG32(True)) if bool(config.get(tirx.instrument_lwp, False)): passes.append(s_tir.transform.InstrumentProfileIntrinsics()) ... if bool(config.get(tirx.use_async_copy, False)): passes.append(s_tir.transform.InjectPTXAsyncCopy()) if bool(config.get(tirx.s_tir.ldg32, False)): passes.append(s_tir.transform.InjectPTXLDG32())这些配置项如tirx.use_async_copy、tirx.s_tir.ldg32、tirx.disable_vectorize、tirx.instrument_bound_checkers、tirx.instrument_lwp均通过tvm.transform.PassContext.current().config读取因此用户可以在构建时通过PassContext(config{...})开启/关闭对应 Pass。管线最终通过tir_pipeline.PIPELINE_MAP[s_tir] default_s_tir_pipeline注册到 TIR 后端的编译管线映射中。如何在自己的编译流程中使用这些 Passtvm.s_tir.transform中的每个工厂函数都返回一个标准tvm.transform.Pass因此可以自由组合到自定义编译流程中。基本用法如下import tvm from tvm import s_tir # 单个 Pass把带 block 的 S-TIR 降级为不可调度的 TIR lower_pass s_tir.transform.LowerOpaqueBlock() # 组合多条 Pass 按顺序执行 pipeline tvm.ir.transform.Sequential([ s_tir.transform.CanonicalizeLoop(), s_tir.transform.ConvertBlocksToOpaque(), s_tir.transform.CompactBufferAllocation(), s_tir.transform.LowerMatchBuffer(), s_tir.transform.LowerOpaqueBlock(), s_tir.transform.StmtSimplify(), ]) mod pipeline(mod) # mod 为 IRModule带参数的 Pass 直接传参例如s_tir.transform.CompactBufferAllocation(is_strictFalse)、s_tir.transform.InjectPTXLDG32(True)、s_tir.transform.ThreadSync(shared)、s_tir.transform.HoistIfThenElse(variantbasic)。需要精确控制调优阶段行为时可使用s_tir.transform.RemoveWeightLayoutRewriteBlock(skip_tensor_rewriteTrue)。如果希望复用仓库默认的完整编排则不需要手工拼装直接使用 python/tvm/s_tir/pipeline.py 提供的default_s_tir_pipeline()并通过PassContext中的配置项如tirx.use_async_copy、tirx.s_tir.ldg32按需启用其中条件性的 Pass。小结tvm.s_tir.transform是 S-TIR 从“可调度 IR”走向“最终代码”之间的核心变换集合。它覆盖了循环规范化、block 降级、buffer 分配收缩、线程绑定与同步、软件流水线与双缓冲、TensorCore/异步拷贝特化、表达式提升、目标校验与性能剖析等完整链路。理解每个 Pass 的职责与顺序既能帮助使用者读懂 python/tvm/s_tir/pipeline.py 中默认管线的每一步意图也能支撑其按需裁剪、组合出满足特定硬件与算子需求的自定义编译流程。建议读者在掌握本文 API 语义的基础上对照 src/s_tir/transform 下同名.cc文件阅读实现以获得对 S-TIR 后端降级最完整的认识。赞分享模型编译深度学习推理引擎【免费下载链接】tvmOpen Machine Learning Compiler Framework项目地址https://gitcode.com/gh_mirrors/tv/tvm点击查看免费下载相关推荐TVM TIRx 编译器变换Compiler Transforms全指南tvm.tirx.transform 的 Pass 体系与降级管线TVM TIRx 编译器变换Compiler Transforms全指南tvm.tirx.transform 的 Pass 体系与降级管线 导读 本文是模型编译深度学习推理引擎Apache TVM TensorIR 深度解析从张量程序抽象、TVMScript 编写到 DLight 与 MetaSchedule 自动化调度Apache TVM TensorIR 深度解析从张量程序抽象、TVMScript 编写到 DLight 与 MetaSchedule 自动化调度 Tenso模型编译深度学习推理引擎TVM Pass Infrastructure 深度解析从 PassContext 到 Pass Instrument 的统一优化管线框架TVM Pass Infrastructure 深度解析从 PassContext 到 Pass Instrument 的统一优化管线框架 本文是 Apach模型编译深度学习推理引擎创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
