- 模型编译
- 深度学习
- 推理引擎
【免费下载链接】tvm
Open Machine Learning Compiler Framework
S-TIR（Schedulable 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.py：
tvm_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）逐步消除或规范化，是整个降级过程的前置步骤。
CanonicalizeLoop
def CanonicalizeLoop(): """Canonicalize the loop to start from zero and use trivial step""" return _ffi_api.CanonicalizeLoop()将循环规范化为从 0 起始、步长为 1（trivial step）的标准形式，为后续 loop partition、向量化等 Pass 提供统一的循环结构。C++ 实现见 src/s_tir/transform/canonicalize_loop.cc。
StmtSimplify
def 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。
ConvertSSA
def 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 复制出一个行为等价的函数。
RenormalizeSplitPattern
def 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。
PlanAndUpdateBufferAllocationLocation
def 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。
ConvertBlocksToOpaque
def 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。
CompactBufferAllocation
def 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 与 LowerOpaqueBlock
def 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.cc；LowerOpaqueBlock：直接删除 opaque block，把 block 内的语句展开到外层作用域。执行完这个 Pass 后，TIR 中不再有任何 block，也就“不可能再被调度”，彻底完成了可调度性到普通 TIR 的过渡。实现见 src/s_tir/transform/lower_opaque_block.cc。
线程绑定、并行与流水线 Pass
这一组面向多线程/多核后端（尤其是 GPU），处理线程绑定、跨线程归约、共享内存同步、软件流水线、双缓冲等。
LiftThreadBinding 与 UnifyThreadBinding
def 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.cc；UnifyThreadBinding：统一blockIdx.x/y/z、threadIdx.x/y/z、vthread.x/y/z的绑定，实现见 src/s_tir/transform/unify_thread_binding.cc。
LowerCrossThreadReduction 与 LowerThreadAllreduce
def 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.cc；LowerThreadAllreduce：降低跨线程 allreduce，见 src/s_tir/transform/lower_thread_allreduce.cc。
两者都服务于 GPU 上多线程协同归约的场景，最终配合ThreadSync插入的同步保证正确性。
ThreadSync
def 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 与 InjectDoubleBuffer
def 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.cc；InjectDoubleBuffer：注入双缓冲语句，并配套注册了InjectDoubleBufferConfig配置对象，实现见 src/s_tir/transform/inject_double_buffer.cc。
InjectSoftwarePipeline 与 ManifestSharedMemoryLocalStage
def 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.cc；ManifestSharedMemoryLocalStage：为 GPU 共享内存访问显式添加 local stage，实现见 src/s_tir/transform/manifest_shared_memory_local_stage.cc。
LowerInitBlock、DefaultGPUSchedule 与 AnnotateIrregularLoop
def 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.cc；DefaultGPUSchedule：为 GPU PrimFunc 设置默认线程绑定，实现见 src/s_tir/transform/default_gpu_schedule.cc；AnnotateIrregularLoop：为不规则循环（如边界不齐的循环）打上注解标记，便于后端做特殊处理，实现见 src/s_tir/transform/annotate_irregular_loop.cc。
GPU 特化：TensorCore、Async Copy 与内存合并
InferFragment、TransformMmaBufferLayout 与 InjectPermutedLayout
def 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.cc；TransformMmaBufferLayout：变换 MMA 相关的 buffer 布局，见 src/s_tir/transform/transform_mma_buffer_layout.cc；InjectPermutedLayout：在 MMA 中注入置换后的布局，见 src/s_tir/transform/inject_permuted_layout.cc。
InjectPTXAsyncCopy 与 LowerAsyncDMA
def 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.cc；InjectPTXAsyncCopy：把 CUDA 上 global 到 shared 的拷贝改写为异步拷贝（cp.async类指令），见 src/s_tir/transform/inject_ptx_async_copy.cc。
InjectPTXLDG32
def InjectPTXLDG32(enable_inject_ptx_intrin=True): """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配置开启时），一次不带参数。
MergeSharedMemoryAllocations
def MergeSharedMemoryAllocations(): """This pass merges multiple TIR-level shared memory allocations into one allocation."""把多个 TIR 层面的共享内存分配合并为一次分配，减少__shared__声明的数量与对齐开销，实现见 src/s_tir/transform/merge_shared_memory_allocations.cc。
表达式提升与代码质量优化 Pass
HoistIfThenElse 与 HoistExpression
def HoistIfThenElse(variant=None): """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"（基础变体）与默认变体两种实现，传入其他值会抛出ValueError；HoistExpression：提升循环不变表达式，实现位于 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 与 UseAssumeToReduceBranches
def 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.cc；RemoveStoreUndef：删除对未定义值的 store，见 src/s_tir/transform/remove_store_undef.cc；UseAssumeToReduceBranches：通过为 padding 区域“过度计算”值来消除 layout 特有的 padding 分支（配合 assume 使用），见 src/s_tir/transform/using_assume_to_reduce_branches.cc。
DecorateDeviceScope
def DecorateDeviceScope(): """Decorate all the function's body as device function."""把函数体整体标记为 device function，见 src/s_tir/transform/decorate_device_scope.cc。
目标感知与诊断 Pass
VerifyVTCMLimit 与 LowerVtcmAlloc
def VerifyVTCMLimit(default_target=None): """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 与 InstrumentProfileIntrinsics
def InstrumentBoundCheckers(): """Instruments bound checkers.""" def InstrumentProfileIntrinsics(): """Insert intrinsic calls to instrument function and loop level profiling."""InstrumentBoundCheckers：插入边界检查器，实现见 src/s_tir/transform/bound_checker.cc；InstrumentProfileIntrinsics：插入 intrinsic 调用以对函数与循环级别做性能剖析，见 src/s_tir/transform/profile_instrumentation.cc。
RemoveWeightLayoutRewriteBlock
def RemoveWeightLayoutRewriteBlock(skip_tensor_rewrite=False): """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 重写 block；skip_tensor_rewrite=True时跳过根据 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 后端的编译管线映射中。
如何在自己的编译流程中使用这些 Pass
tvm.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_strict=False)、s_tir.transform.InjectPTXLDG32(True)、s_tir.transform.ThreadSync("shared")、s_tir.transform.HoistIfThenElse(variant="basic")。需要精确控制调优阶段行为时，可使用s_tir.transform.RemoveWeightLayoutRewriteBlock(skip_tensor_rewrite=True)。
如果希望复用仓库默认的完整编排，则不需要手工拼装，直接使用 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 后端降级最完整的认识。
- 模型编译
- 深度学习
- 推理引擎
【免费下载链接】tvm
Open Machine Learning Compiler Framework
相关推荐
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），仅供参考
