From 18159f315839494bedb317735854768ccb31d6ea Mon Sep 17 00:00:00 2001
From: tlopex <820958424@qq.com>
Date: Sat, 15 Aug 2026 00:56:49 -0400
Subject: [PATCH 1/4] Update TIRx compiler internals guide

---
 appendix/index.md                        |   2 +-
 index.md                                 |   2 +-
 tirx_guide/arch/index.rst                |  10 +-
 tirx_guide/arch/lowering_pipeline.rst    | 133 +++++++++++++++--------
 zh/appendix/index.md                     |   2 +-
 zh/index.md                              |   2 +-
 zh/tirx_guide/arch/index.rst             |   9 +-
 zh/tirx_guide/arch/lowering_pipeline.rst | 131 ++++++++++++++--------
 8 files changed, 180 insertions(+), 111 deletions(-)

diff --git a/appendix/index.md b/appendix/index.md
index d836e0bd..c0b27807 100644
--- a/appendix/index.md
+++ b/appendix/index.md
@@ -6,8 +6,8 @@ The main text runs through Parts I–IV. The Reference section collects material
 | Need | Where |
 |------|-----|
 | Look up a TIRx language feature | **{ref}`chap_language_reference`** |
-| Debug asynchronous GEMM/FA hangs, crashes, wrong results, or slowdowns | **{ref}`chap_warp_spec_debug`** |
 | Compiler internals (the lowering pipeline) | **{ref}`chap_arch`** |
+| Debug asynchronous GEMM/FA hangs, crashes, wrong results, or slowdowns | **{ref}`chap_warp_spec_debug`** |
 
 For the complete `tvm.tirx` Python API, see the
 [upstream TVM documentation](https://tvm.apache.org/docs/).
diff --git a/index.md b/index.md
index 70b264a8..5e903ef6 100644
--- a/index.md
+++ b/index.md
@@ -87,6 +87,6 @@ chapter_flash_attention/index
 
 appendix/index
 tirx_guide/language_reference/index
+TIRx Compiler Internals <tirx_guide/arch/lowering_pipeline>
 appendix/debugging_warp_specialized
-tirx_guide/arch/index
 ```
diff --git a/tirx_guide/arch/index.rst b/tirx_guide/arch/index.rst
index 12533ef1..be24d4cb 100644
--- a/tirx_guide/arch/index.rst
+++ b/tirx_guide/arch/index.rst
@@ -15,14 +15,10 @@
     specific language governing permissions and limitations
     under the License.
 
-.. _chap_arch:
+:orphan:
 
 Compiler Internals
 ==================
 
-Internals of the TIRx compiler, for contributors.
-
-.. toctree::
-   :maxdepth: 1
-
-   lowering_pipeline
+The compiler-internals guide is now available at
+:ref:`TIRx Compiler Internals: Compilation and Lowering Pipeline <chap_arch>`.
diff --git a/tirx_guide/arch/lowering_pipeline.rst b/tirx_guide/arch/lowering_pipeline.rst
index 1fc17a63..f056c09c 100644
--- a/tirx_guide/arch/lowering_pipeline.rst
+++ b/tirx_guide/arch/lowering_pipeline.rst
@@ -15,19 +15,24 @@
     specific language governing permissions and limitations
     under the License.
 
-TIRx lowering pipeline
-======================
+.. _chap_arch:
 
-``tvm.compile(mod, target, tir_pipeline="tirx")`` passes a TIRx module through the
-**tirx pipeline**, an ordered sequence of TIR passes. These passes lower the
-high-level constructs in the source—tile primitives, ``TileLayout``-typed buffers,
-and execution-scope IDs—and split the module into **host** and **device** functions.
-The CUDA backend then generates source for the device functions. The pipeline is defined in
-``python/tvm/tirx/compilation_pipeline.py`` (``tirx_pipeline``); this page walks the
-passes in order.
+TIRx Compiler Internals: Compilation and Lowering Pipeline
+===========================================================
 
-Where it sits
--------------
+``tvm.compile(mod, target, tir_pipeline="tirx")`` turns an authored TIRx module
+into host launcher code and device code. The work does not happen in one step:
+TIRx-specific constructs are lowered first, general-purpose TIRx normalization
+and legalization passes then process the result, and the module is finally split
+and prepared for code generation.
+
+The exact sequence is defined in `compilation_pipeline.py
+<https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/compilation_pipeline.py>`_.
+This page explains where that sequence sits in ``tvm.compile``, what each stage
+changes, and where the host and device paths diverge.
+
+The overall compilation path
+----------------------------
 
 ``tvm.compile`` first binds the target, runs the **tirx pipeline** (the module-level
 passes below), then applies **finalization** passes separately to the host and
@@ -40,112 +45,146 @@ generator:
                                           │
                                           └──────────▶  device func ──device finalize──▶  CUDA
 
-The passes
-----------
+Pass order inside ``tirx_pipeline``
+-----------------------------------
 
-The ``tirx_pipeline`` module pass applies this exact sequence (a few are gated by
-``PassContext`` config):
+The pipeline is organized into the 19 steps below. Common-subexpression
+elimination is optional, while vectorization and unrolling behavior can be
+controlled through ``PassContext``:
 
 .. list-table::
    :header-rows: 1
-   :widths: 6 32 62
+   :widths: 6 24 24 46
 
    * - #
+     - Stage
      - Pass
      - What it does
    * - 1
+     - TIRx lowering
      - ``LowerTIRx``
      - the core lowering — see `Inside LowerTIRx`_ below
    * - 2
+     - TIR normalization
      - ``UnifyThreadBinding``
      - merges equivalent thread-axis bindings so each ``threadIdx`` / ``blockIdx``
        axis is declared once
    * - 3
+     - TIR normalization
      - ``StmtSimplify``
      - statement-level arithmetic simplification (the arith analyzer)
    * - 4
+     - TIR normalization
      - ``LowerTIRxOpaque``
-     - lowers remaining opaque TIRx constructs to plain TIR
+     - lowers remaining opaque constructs, including thread-binding loops,
+       unit loops, and pragma annotations
    * - 5
+     - TIR normalization
      - ``FlattenBuffer``
-     - flattens multi-dimensional ``BufferLoad`` / ``BufferStore`` to 1-D
+     - flattens the remaining multi-dimensional TIR ``BufferLoad`` /
+       ``BufferStore`` accesses to 1-D
    * - 6
+     - Compute legalization
      - ``BF16ComputeLegalize``
      - rewrites ``bfloat16`` compute to a legal (f32-up-cast) form
    * - 7
+     - TIR normalization
      - ``NarrowDataType(32)``
      - narrows index/loop scalar ``Expr`` types to 32-bit where provably safe
    * - 8
+     - Loop lowering
      - ``VectorizeLoop``
-     - turns ``T.vectorized`` loops into vector ops (skipped if
-       ``tir.disable_vectorize``)
+     - lowers ``T.vectorized`` loops to vector operations; when
+       ``tir.disable_vectorize`` is set, it instead scalarizes those loops
    * - 9
+     - Loop lowering
      - ``UnrollLoop``
-     - unrolls loops marked ``T.unroll`` (and small constant loops)
+     - unrolls loops marked ``T.unroll``; ordinary constant loops are
+       auto-unrolled only when the corresponding config or pragma enables it
    * - 10
+     - TIR normalization
      - ``StmtSimplify``
      - simplify again, now that vectorize/unroll exposed constants
    * - 11
+     - TIR normalization
      - ``CommonSubexprElim``
      - hoists repeated subexpressions into temporaries (skipped if
        ``tir.disable_cse_tir``)
    * - 12
+     - Compute legalization
      - ``FP8ComputeLegalize``
      - rewrites ``float8`` compute to a legal form
    * - 13
+     - Validation and ABI
      - ``VerifyMemory``
      - checks no host-side code directly dereferences device memory (a safety gate)
    * - 14
+     - Validation and ABI
      - ``AnnotateEntryFunc``
-     - marks the single PrimFunc as the module entry point
+     - marks the sole function, or the sole externally visible PrimFunc in a
+       multi-function module, as the entry point
    * - 15
+     - Validation and ABI
      - ``SplitHostDevice``
-     - splits each kernel into a **host** function and a **device** function at the
-       ``launch_thread`` boundary
+     - identifies device regions, splits host and device PrimFuncs, and lowers
+       host-to-device calls to the kernel-launch ABI
    * - 16
+     - Validation and ABI
      - ``LowerIket``
      - removes frontend-only NVIDIA IKET annotations for normal builds, or emits
        IKET metadata and placeholders when the IRModule is explicitly IKET-enabled
    * - 17
+     - Validation and ABI
      - ``MakePackedAPI``
      - rewrites the host function to the packed-func ABI (the launcher TVM calls)
    * - 18
+     - Storage legalization
      - ``FP8StorageLegalize``
-     - legalizes ``float8`` storage (packing into supported container types)
+     - legalizes ``float8`` storage to ``uint8`` containers
    * - 19
+     - Storage legalization
      - ``BF16StorageLegalize``
-     - legalizes ``bfloat16`` storage
+     - legalizes ``bfloat16`` storage to ``uint16`` containers
+
+Host and device finalization
+----------------------------
 
-**Finalization** then runs per function kind:
+The 19 listed steps form ``tirx_pipeline``. After that module-level pipeline,
+``tvm.compile`` runs a different finalization sequence for each function kind:
 
 - **host**: ``LowerTVMBuiltin`` (lower ``tvm_*`` builtins), ``LowerIntrin``
   (target-specific intrinsics)
 - **device**: ``LowerWarpMemory`` (warp-scoped buffers → shuffles), ``StmtSimplify``,
   ``LowerIntrin``
 
-Inside LowerTIRx
-----------------
+Inside ``LowerTIRx``
+--------------------
 
-``LowerTIRx`` is itself a small sequence (``src/tirx/transform/lower_tirx.cc``):
+In a normal build, ``LowerTIRx`` is itself a two-pass sequence defined in
+`lower_tirx.cc
+<https://github.com/apache/tvm/blob/v0.26.0/src/tirx/transform/lower_tirx.cc>`_:
 
 .. code-block:: text
 
     LowerTIRx = Sequential([ TilePrimitiveDispatch, LowerTIRxCleanup ])
 
 - **``TilePrimitiveDispatch``** selects a backend variant for every
-  ``TilePrimitiveCall`` (``copy``, ``gemm``, ``reduction``, …) and replaces the
-  call with the implementation emitted by that variant.
+  ``TilePrimitiveCall`` (``copy``, ``gemm``, ``reduction``, …), replaces the
+  call with the selected implementation, and resolves execution-scope IDs such
+  as ``T.cta_id`` and ``T.thread_id`` into kernel launch parameters and bindings.
 - **``LowerTIRxCleanup``** runs the ``LayoutApplier``: it resolves every
   ``TileLayout``-typed buffer access into concrete physical address arithmetic
-  (``addr = data + elem_offset + layout.apply(*coord, shape=shape)``), flattens the buffers, and
-  lowers the execution-scope ids (``T.cta_id`` / ``T.thread_id`` / … →
-  ``blockIdx`` / ``threadIdx`` via ``launch_thread``).
+  (``addr = data + elem_offset + layout.apply(*coord, shape=shape)``), replaces
+  layout-aware buffer parameters with physical views, and removes explicit
+  buffer offsets.
 
-So after ``LowerTIRx`` the module is plain TIR: no tile primitives, no
-``TileLayout`` indirection, scope ids resolved to thread axes.
+After ``LowerTIRx``, tile primitives and ``TileLayout`` indirection are gone,
+and execution-scope IDs have been resolved. Some opaque TIRx constructs still
+remain; the later ``LowerTIRxOpaque`` pass converts those before
+``tirx.transform.FlattenBuffer`` flattens ordinary TIR buffer accesses.
 
-A worked example
-----------------
+End-to-end IR evolution
+-----------------------
 
 Take a one-line scale kernel:
 
@@ -158,19 +197,23 @@ Take a one-line scale kernel:
         T.device_entry(); bx = T.cta_id([1]); tx = T.thread_id([256])
         B[tx] = A[tx] * T.float32(2.0)
 
-**After ``LowerTIRx``** the scope ids are real thread axes and the layout is applied
-(``A_1`` / ``B_1`` are the flattened 1-D views):
+This simple 1-D kernel has no nontrivial ``TileLayout``; it chiefly shows how
+``LowerTIRx`` turns scope IDs into real thread axes. The core body looks like the
+following excerpt. Buffer declarations and an unused warp-ID binding are omitted;
+``A_1`` and ``B_1`` are the materialized physical views:
 
 .. code-block:: python
 
+    # match_buffer / decl_buffer declarations omitted
     with T.launch_thread("blockIdx.x", 1) as blockIdx_x:
         threadIdx_x = T.launch_thread("threadIdx.x", 256)
         bx: T.let = blockIdx_x
         tx: T.let = threadIdx_x
         B_1[threadIdx_x] = A_1[threadIdx_x] * T.float32(2.0)
 
-**After ``SplitHostDevice`` + ``LowerIket`` + ``MakePackedAPI``** the one function has become two —
-a host launcher and a device kernel:
+``SplitHostDevice`` then turns the single function into a host launcher and a
+device kernel. ``MakePackedAPI`` later rewrites the host launcher to TVM's
+packed-function ABI:
 
 .. code-block:: text
 
@@ -183,8 +226,8 @@ a host launcher and a device kernel:
 The CUDA backend then renders ``scale_kernel`` to the ``__global__`` function
 (``B_ptr[threadIdx.x] = A_ptr[threadIdx.x] * 2.0f``).
 
-Reproduce it yourself
----------------------
+Inspecting intermediate IR and generated code
+----------------------------------------------
 
 You can run any prefix of the pipeline by hand to inspect a stage — this is how the
 IR snippets across these docs were produced:
diff --git a/zh/appendix/index.md b/zh/appendix/index.md
index f92847cc..0060cbce 100644
--- a/zh/appendix/index.md
+++ b/zh/appendix/index.md
@@ -6,8 +6,8 @@
 | 需要查询的内容 | 对应页面 |
 |---|---|
 | TIRx 语言特性的准确写法和语义 | **{ref}`chap_language_reference`** |
-| 排查异步 GEMM 或 Flash Attention kernel 的卡死、崩溃、错误结果和性能下降 | **{ref}`chap_warp_spec_debug`** |
 | 编译器内部机制与 lowering 流程 | **{ref}`chap_arch`** |
+| 排查异步 GEMM 或 Flash Attention kernel 的卡死、崩溃、错误结果和性能下降 | **{ref}`chap_warp_spec_debug`** |
 
 完整的 `tvm.tirx` Python API 请参阅
 [TVM 官方文档](https://tvm.apache.org/docs/)。
diff --git a/zh/index.md b/zh/index.md
index b5a2eb15..b0f43df4 100644
--- a/zh/index.md
+++ b/zh/index.md
@@ -64,6 +64,6 @@ chapter_flash_attention/index
 
 appendix/index
 tirx_guide/language_reference/index
+TIRx 编译器内部机制 <tirx_guide/arch/lowering_pipeline>
 appendix/debugging_warp_specialized
-tirx_guide/arch/index
 ```
diff --git a/zh/tirx_guide/arch/index.rst b/zh/tirx_guide/arch/index.rst
index a4ab2465..0c0ed23b 100644
--- a/zh/tirx_guide/arch/index.rst
+++ b/zh/tirx_guide/arch/index.rst
@@ -15,14 +15,9 @@
     specific language governing permissions and limitations
     under the License.
 
-.. _chap_arch:
+:orphan:
 
 编译器内部机制
 ==============
 
-本节面向 TIRx 贡献者，介绍编译器内部的工作方式。
-
-.. toctree::
-   :maxdepth: 1
-
-   lowering_pipeline
+完整内容已移至 :ref:`TIRx 编译器内部机制：编译与 Lowering 流水线 <chap_arch>`。
diff --git a/zh/tirx_guide/arch/lowering_pipeline.rst b/zh/tirx_guide/arch/lowering_pipeline.rst
index 16800e47..af60d211 100644
--- a/zh/tirx_guide/arch/lowering_pipeline.rst
+++ b/zh/tirx_guide/arch/lowering_pipeline.rst
@@ -15,20 +15,23 @@
     specific language governing permissions and limitations
     under the License.
 
-TIRx Lowering Pipeline
-======================
+.. _chap_arch:
 
-调用 ``tvm.compile(mod, target, tir_pipeline="tirx")`` 时，编译器会将输入的
-TIRx module 依次送入一组 TIR passes，这组 passes 称为 **tirx pipeline**。
-它负责把 tile primitives、使用 ``TileLayout`` 的 buffers 和 execution-scope
-ids 等高层结构逐步转换为彼此分离的 **host** 与 **device** functions，最后再
-由 CUDA backend 生成源码。
+TIRx 编译器内部机制：编译与 Lowering 流水线
+================================================
 
-Pipeline 定义在 ``python/tvm/tirx/compilation_pipeline.py`` 的
-``tirx_pipeline`` 中。下面按执行顺序介绍其中的 passes。
+``tvm.compile(mod, target, tir_pipeline="tirx")`` 会把编写好的 TIRx module
+转换成 host launcher 和 device code。这个过程并非一次完成：编译器先处理
+TIRx 特有的结构，再用通用的 TIRx 规范化与合法化 passes 处理结果，最后拆分
+module 并交给后端生成代码。
 
-Pipeline 在编译流程中的位置
----------------------------
+完整顺序定义在 `compilation_pipeline.py
+<https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/compilation_pipeline.py>`_。
+本页先说明这条流水线在 ``tvm.compile`` 中的位置，再按阶段解释各个 pass
+改变了什么，以及 host 与 device 两条路径从哪里分开。
+
+整体编译路径
+------------
 
 ``tvm.compile`` 首先绑定 target，再运行下面的 module-level **tirx
 pipeline**。随后，host 和 device functions 分别经过 finalization passes，
@@ -40,113 +43,142 @@ device function 最终交给 CUDA code generator：
                                           │
                                           └──────────▶  device func ──device finalize──▶  CUDA
 
-Pass 执行顺序
--------------
+``tirx_pipeline`` 的 Pass 执行顺序
+-----------------------------------
 
-``tirx_pipeline`` 依次执行下表中的 passes，其中少数 pass 会由
-``PassContext`` 配置控制是否启用：
+``tirx_pipeline`` 按下表中的 19 个步骤组织。公共子表达式消除是可选项，
+vectorization 和 unrolling 的行为也可以通过 ``PassContext`` 控制：
 
 .. list-table::
    :header-rows: 1
-   :widths: 6 32 62
+   :widths: 6 24 24 46
 
    * - #
+     - 阶段
      - Pass
      - 作用
    * - 1
+     - TIRx lowering
      - ``LowerTIRx``
      - 完成 TIRx 的核心转换，详见下方 `LowerTIRx 内部做了什么`_
    * - 2
+     - TIR 规范化
      - ``UnifyThreadBinding``
      - 合并等价的 thread-axis bindings，使每个 ``threadIdx`` / ``blockIdx``
        axis 只声明一次
    * - 3
+     - TIR 规范化
      - ``StmtSimplify``
      - 使用 arithmetic analyzer 简化 statement 中的算术表达式
    * - 4
+     - TIR 规范化
      - ``LowerTIRxOpaque``
-     - 将剩余的 opaque TIRx constructs 转换为普通 TIR
+     - 处理剩余的 opaque constructs，包括 thread-binding loops、unit loops
+       和 pragma annotations
    * - 5
+     - TIR 规范化
      - ``FlattenBuffer``
-     - 将多维 ``BufferLoad`` / ``BufferStore`` 展平为一维访问
+     - 将剩余的多维 TIR ``BufferLoad`` / ``BufferStore`` 展平为一维访问
    * - 6
+     - 计算合法化
      - ``BF16ComputeLegalize``
      - 将 ``bfloat16`` 计算改写为合法形式，其中计算会提升到 f32
    * - 7
+     - TIR 规范化
      - ``NarrowDataType(32)``
      - 在能够证明安全时，将 index 和 loop 的 scalar ``Expr`` type 缩窄为 32 bits
    * - 8
+     - Loop lowering
      - ``VectorizeLoop``
-     - 将 ``T.vectorized`` loops 改写为 vector operations；设置
-       ``tir.disable_vectorize`` 时跳过
+     - 将 ``T.vectorized`` loops 转换为 vector operations；设置
+       ``tir.disable_vectorize`` 时，改为将这些 loops scalarize
    * - 9
+     - Loop lowering
      - ``UnrollLoop``
-     - 展开标记为 ``T.unroll`` 的 loops，以及较小的常量 loops
+     - 展开标记为 ``T.unroll`` 的 loops；普通常量 loops 只有在相应配置或
+       pragma 启用时才会自动展开
    * - 10
+     - TIR 规范化
      - ``StmtSimplify``
      - Vectorize 和 unroll 暴露出更多常量后，再次执行简化
    * - 11
+     - TIR 规范化
      - ``CommonSubexprElim``
      - 将重复的子表达式提取为临时变量；设置 ``tir.disable_cse_tir`` 时跳过
    * - 12
+     - 计算合法化
      - ``FP8ComputeLegalize``
      - 将 ``float8`` 计算改写为合法形式
    * - 13
+     - 校验与 ABI
      - ``VerifyMemory``
      - 检查 host 代码没有直接解引用 device memory
    * - 14
+     - 校验与 ABI
      - ``AnnotateEntryFunc``
-     - 将 module 中唯一的 PrimFunc 标记为入口函数
+     - 将唯一 function 标记为入口；对于多 function module，则标记其中唯一
+       对外可见的 PrimFunc
    * - 15
+     - 校验与 ABI
      - ``SplitHostDevice``
-     - 在 ``launch_thread`` 边界处，将每个 kernel 拆分为 **host** function
-       和 **device** function
+     - 识别 device regions，拆分 host 与 device PrimFuncs，并将 host 侧调用
+       转换为 kernel-launch ABI
    * - 16
+     - 校验与 ABI
      - ``LowerIket``
      - 普通 build 中移除 frontend-only NVIDIA IKET annotations；IRModule 显式
        启用 IKET 时则生成 IKET metadata 和 placeholders
    * - 17
+     - 校验与 ABI
      - ``MakePackedAPI``
      - 将 host function 改写为 TVM launcher 使用的 packed-function ABI
    * - 18
+     - Storage 合法化
      - ``FP8StorageLegalize``
-     - 将 ``float8`` storage 打包为 backend 支持的容器类型
+     - 将 ``float8`` storage 转换为 ``uint8`` container
    * - 19
+     - Storage 合法化
      - ``BF16StorageLegalize``
-     - 将 ``bfloat16`` storage 改写为合法形式
+     - 将 ``bfloat16`` storage 转换为 ``uint16`` container
 
-之后，编译器会根据 function 类型分别执行 **finalization**：
+Host 与 Device Finalization
+---------------------------
+
+上面列出的 19 个步骤组成 ``tirx_pipeline``。这条 module-level pipeline
+结束后，``tvm.compile`` 会根据 function 类型分别执行 finalization：
 
 - **host**：``LowerTVMBuiltin`` 处理 ``tvm_*`` builtins，``LowerIntrin``
   处理 target-specific intrinsics。
 - **device**：``LowerWarpMemory`` 将 warp-scoped buffers 转换为
   shuffles，随后执行 ``StmtSimplify`` 和 ``LowerIntrin``。
 
-LowerTIRx 内部做了什么
------------------------
+``LowerTIRx`` 内部做了什么
+---------------------------
 
-``LowerTIRx`` 本身由两个 passes 组成，定义在
-``src/tirx/transform/lower_tirx.cc``：
+正常编译时，``LowerTIRx`` 本身由两个 passes 组成，定义在
+`lower_tirx.cc
+<https://github.com/apache/tvm/blob/v0.26.0/src/tirx/transform/lower_tirx.cc>`_：
 
 .. code-block:: text
 
     LowerTIRx = Sequential([ TilePrimitiveDispatch, LowerTIRxCleanup ])
 
-- **``TilePrimitiveDispatch``** 根据选中的 backend dispatch，将每个
-  ``TilePrimitiveCall``（``copy``、``gemm``、``reduction`` 等）替换为对应
-  的实现。
+- **``TilePrimitiveDispatch``** 根据 backend dispatch，为每个
+  ``TilePrimitiveCall``（``copy``、``gemm``、``reduction`` 等）选择具体实现，
+  同时把 ``T.cta_id``、``T.thread_id`` 等 execution-scope IDs 解析为 kernel
+  launch parameters 和对应的 bindings。
 - **``LowerTIRxCleanup``** 运行 ``LayoutApplier``，将使用
   ``TileLayout`` 的 buffer access 变成具体的物理地址计算
-  （``addr = data + elem_offset + layout.apply(*coord, shape=shape)``），再展平 buffers，
-  并将 execution-scope ids 转换为 thread axes，例如
-  ``T.cta_id`` / ``T.thread_id`` 通过 ``launch_thread`` 变为
-  ``blockIdx`` / ``threadIdx``。
+  （``addr = data + elem_offset + layout.apply(*coord, shape=shape)``），把带 layout
+  的 buffer parameters 替换成物理 views，并移除显式的 buffer offsets。
 
-完成 ``LowerTIRx`` 后，module 中只剩普通 TIR：tile primitives 已经展开，
-``TileLayout`` 间接层已经消失，scope ids 也已经解析为 thread axes。
+完成 ``LowerTIRx`` 后，tile primitives 和 ``TileLayout`` 间接层已经消失，
+execution-scope IDs 也已经解析。此时仍有少量 opaque TIRx constructs；后续的
+``LowerTIRxOpaque`` 会先处理这些结构，后续的
+``tirx.transform.FlattenBuffer`` pass 再展平普通 TIR 中的 buffer access。
 
-完整示例
---------
+端到端 IR 演化示例
+------------------
 
 以下面的 scale kernel 为例：
 
@@ -159,19 +191,22 @@ LowerTIRx 内部做了什么
         T.device_entry(); bx = T.cta_id([1]); tx = T.thread_id([256])
         B[tx] = A[tx] * T.float32(2.0)
 
-执行 ``LowerTIRx`` 后，scope ids 已经变成真实的 thread axes，layout 也已经
-应用到 buffer access 上。这里的 ``A_1`` 和 ``B_1`` 是展平后的一维 views：
+这个简单的一维 kernel 没有非平凡的 ``TileLayout``，主要用来展示
+``LowerTIRx`` 如何将 scope IDs 转换成真实的 thread axes。下面只摘录核心
+body，省略 buffer declarations 和未使用的 warp-ID binding；``A_1`` 与
+``B_1`` 是生成的物理 views：
 
 .. code-block:: python
 
+    # 省略 match_buffer / decl_buffer declarations
     with T.launch_thread("blockIdx.x", 1) as blockIdx_x:
         threadIdx_x = T.launch_thread("threadIdx.x", 256)
         bx: T.let = blockIdx_x
         tx: T.let = threadIdx_x
         B_1[threadIdx_x] = A_1[threadIdx_x] * T.float32(2.0)
 
-经过 ``SplitHostDevice``、``LowerIket`` 和 ``MakePackedAPI`` 后，一个 function 会拆成 host
-launcher 和 device kernel：
+``SplitHostDevice`` 随后将单个 function 拆成 host launcher 和 device kernel，
+``MakePackedAPI`` 再将 host launcher 转换为 TVM 的 packed-function ABI：
 
 .. code-block:: text
 
@@ -184,8 +219,8 @@ launcher 和 device kernel：
 CUDA backend 随后将 ``scale_kernel`` 生成 ``__global__`` function：
 ``B_ptr[threadIdx.x] = A_ptr[threadIdx.x] * 2.0f``。
 
-手动检查中间结果
-----------------
+检查中间 IR 与生成代码
+----------------------
 
 可以手动运行 pipeline 的任意前缀，检查某个阶段的 IR。本书中的 IR 片段也是
 用这种方式生成的：

From 6f28eaa271293e808a74bb8871dbdee06d2cd4b7 Mon Sep 17 00:00:00 2001
From: tlopex <820958424@qq.com>
Date: Sat, 15 Aug 2026 03:05:15 -0400
Subject: [PATCH 2/4] Update TIRx compiler internals guide

---
 tirx_guide/arch/index.rst                |   3 +-
 tirx_guide/arch/lowering_pipeline.rst    | 242 ++++++++++++++---------
 zh/tirx_guide/arch/index.rst             |   2 +-
 zh/tirx_guide/arch/lowering_pipeline.rst | 220 +++++++++++----------
 4 files changed, 277 insertions(+), 190 deletions(-)

diff --git a/tirx_guide/arch/index.rst b/tirx_guide/arch/index.rst
index be24d4cb..5a213844 100644
--- a/tirx_guide/arch/index.rst
+++ b/tirx_guide/arch/index.rst
@@ -20,5 +20,4 @@
 Compiler Internals
 ==================
 
-The compiler-internals guide is now available at
-:ref:`TIRx Compiler Internals: Compilation and Lowering Pipeline <chap_arch>`.
+The compiler-internals guide is now available at :ref:`chap_arch`.
diff --git a/tirx_guide/arch/lowering_pipeline.rst b/tirx_guide/arch/lowering_pipeline.rst
index f056c09c..792fdcbd 100644
--- a/tirx_guide/arch/lowering_pipeline.rst
+++ b/tirx_guide/arch/lowering_pipeline.rst
@@ -17,47 +17,59 @@
 
 .. _chap_arch:
 
-TIRx Compiler Internals: Compilation and Lowering Pipeline
-===========================================================
+TIRx Compiler Internals
+=======================
 
-``tvm.compile(mod, target, tir_pipeline="tirx")`` turns an authored TIRx module
-into host launcher code and device code. The work does not happen in one step:
-TIRx-specific constructs are lowered first, general-purpose TIRx normalization
-and legalization passes then process the result, and the module is finally split
-and prepared for code generation.
+``tvm.compile(mod, target, tir_pipeline="tirx")`` takes a TIRx module and
+eventually produces two pieces of code: a CPU-side launcher that prepares the
+arguments and launches a GPU kernel, and the GPU kernel that performs the
+computation. The compiler reaches that result through an ordered series of
+passes. Each pass performs a particular transformation, validation, or
+annotation on the IR.
 
-The exact sequence is defined in `compilation_pipeline.py
+The exact pass order is defined in Apache TVM's `compilation_pipeline.py
 <https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/compilation_pipeline.py>`_.
-This page explains where that sequence sits in ``tvm.compile``, what each stage
-changes, and where the host and device paths diverge.
 
 The overall compilation path
 ----------------------------
 
-``tvm.compile`` first binds the target, runs the **tirx pipeline** (the module-level
-passes below), then applies **finalization** passes separately to the host and
-device functions, and finally hands each device function to the CUDA code
-generator:
+The ``target`` identifies the hardware and code-generation backend. The example
+below uses CUDA for the device and LLVM for the host. ``tvm.compile`` first
+attaches that target information to the module and then runs the module-level
+**tirx pipeline**. Once the pipeline has separated the CPU-side host function
+from the GPU-side device function, each follows a target-specific finalization
+path before code generation:
 
 .. code-block:: text
 
-    authored TIRx  ──BindTarget──▶  tirx_pipeline  ──▶  host func  ──host finalize──▶  C/LLVM
-                                          │
-                                          └──────────▶  device func ──device finalize──▶  CUDA
+    authored TIRx
+          │ BindTarget
+          ▼
+    tirx_pipeline
+    (SplitHostDevice creates the two paths)
+          ├── host PrimFunc   ──host finalization──▶ C/LLVM
+          └── device PrimFunc ─device finalization─▶ CUDA
+
+A ``PrimFunc`` is TIR's representation of a function. The host PrimFunc above is
+the CPU-side launcher, while the device PrimFunc is the GPU kernel.
+``Finalization`` refers to the last target-specific transformations performed
+before code generation.
 
 Pass order inside ``tirx_pipeline``
 -----------------------------------
 
-The pipeline is organized into the 19 steps below. Common-subexpression
-elimination is optional, while vectorization and unrolling behavior can be
-controlled through ``PassContext``:
+The table lists the 19 pipeline steps in execution order. An ABI is the calling
+convention between functions; the ABI passes below adapt ordinary TIR functions
+to forms that the runtime can invoke. ``PassContext`` holds compiler options:
+common-subexpression elimination can be disabled, and it also controls aspects
+of vectorization and unrolling.
 
 .. list-table::
    :header-rows: 1
    :widths: 6 24 24 46
 
    * - #
-     - Stage
+     - Category
      - Pass
      - What it does
    * - 1
@@ -72,12 +84,12 @@ controlled through ``PassContext``:
    * - 3
      - TIR normalization
      - ``StmtSimplify``
-     - statement-level arithmetic simplification (the arith analyzer)
+     - simplifies arithmetic expressions in the IR
    * - 4
      - TIR normalization
      - ``LowerTIRxOpaque``
-     - lowers remaining opaque constructs, including thread-binding loops,
-       unit loops, and pragma annotations
+     - converts thread-binding loops, eliminates unannotated unit loops, and
+       normalizes loop pragmas
    * - 5
      - TIR normalization
      - ``FlattenBuffer``
@@ -86,11 +98,12 @@ controlled through ``PassContext``:
    * - 6
      - Compute legalization
      - ``BF16ComputeLegalize``
-     - rewrites ``bfloat16`` compute to a legal (f32-up-cast) form
+     - when the target lacks native ``bfloat16`` compute, promotes operations to
+       ``float32`` and rewrites them into a legal form
    * - 7
      - TIR normalization
      - ``NarrowDataType(32)``
-     - narrows index/loop scalar ``Expr`` types to 32-bit where provably safe
+     - narrows index expressions and loop variables to 32 bits where provably safe
    * - 8
      - Loop lowering
      - ``VectorizeLoop``
@@ -104,7 +117,7 @@ controlled through ``PassContext``:
    * - 10
      - TIR normalization
      - ``StmtSimplify``
-     - simplify again, now that vectorize/unroll exposed constants
+     - simplifies again after vectorization and unrolling expose more constants
    * - 11
      - TIR normalization
      - ``CommonSubexprElim``
@@ -113,11 +126,12 @@ controlled through ``PassContext``:
    * - 12
      - Compute legalization
      - ``FP8ComputeLegalize``
-     - rewrites ``float8`` compute to a legal form
+     - when the target lacks native ``float8`` compute, promotes operations to a
+       supported type (``float32`` by default)
    * - 13
      - Validation and ABI
      - ``VerifyMemory``
-     - checks no host-side code directly dereferences device memory (a safety gate)
+     - ensures that host-side code does not directly dereference device memory
    * - 14
      - Validation and ABI
      - ``AnnotateEntryFunc``
@@ -131,20 +145,20 @@ controlled through ``PassContext``:
    * - 16
      - Validation and ABI
      - ``LowerIket``
-     - removes frontend-only NVIDIA IKET annotations for normal builds, or emits
-       IKET metadata and placeholders when the IRModule is explicitly IKET-enabled
+     - removes NVIDIA IKET annotations in normal builds, or lowers them for
+       tracing when IKET is enabled
    * - 17
      - Validation and ABI
      - ``MakePackedAPI``
-     - rewrites the host function to the packed-func ABI (the launcher TVM calls)
+     - rewrites the host function to the packed-function ABI used by the TVM runtime
    * - 18
      - Storage legalization
      - ``FP8StorageLegalize``
-     - legalizes ``float8`` storage to ``uint8`` containers
+     - when the target lacks native ``float8`` storage, uses ``uint8`` containers
    * - 19
      - Storage legalization
      - ``BF16StorageLegalize``
-     - legalizes ``bfloat16`` storage to ``uint16`` containers
+     - when the target lacks native ``bfloat16`` storage, uses ``uint16`` containers
 
 Host and device finalization
 ----------------------------
@@ -152,15 +166,17 @@ Host and device finalization
 The 19 listed steps form ``tirx_pipeline``. After that module-level pipeline,
 ``tvm.compile`` runs a different finalization sequence for each function kind:
 
-- **host**: ``LowerTVMBuiltin`` (lower ``tvm_*`` builtins), ``LowerIntrin``
-  (target-specific intrinsics)
-- **device**: ``LowerWarpMemory`` (warp-scoped buffers → shuffles), ``StmtSimplify``,
-  ``LowerIntrin``
+- **host**: ``LowerTVMBuiltin`` (lowers ``tvm_*`` builtins), ``LowerIntrin``
+  (lowers target-specific intrinsics)
+- **device**: ``LowerWarpMemory`` (lowers warp-scoped buffers to shuffles),
+  ``StmtSimplify``, ``LowerIntrin``
 
 Inside ``LowerTIRx``
 --------------------
 
-In a normal build, ``LowerTIRx`` is itself a two-pass sequence defined in
+``LowerTIRx`` has two main jobs: choosing concrete implementations for tile-level
+operations, and turning logical data layouts into physical memory indices. Its
+core transformation is the following two-pass sequence, defined in Apache TVM's
 `lower_tirx.cc
 <https://github.com/apache/tvm/blob/v0.26.0/src/tirx/transform/lower_tirx.cc>`_:
 
@@ -168,82 +184,132 @@ In a normal build, ``LowerTIRx`` is itself a two-pass sequence defined in
 
     LowerTIRx = Sequential([ TilePrimitiveDispatch, LowerTIRxCleanup ])
 
-- **``TilePrimitiveDispatch``** selects a backend variant for every
-  ``TilePrimitiveCall`` (``copy``, ``gemm``, ``reduction``, …), replaces the
-  call with the selected implementation, and resolves execution-scope IDs such
-  as ``T.cta_id`` and ``T.thread_id`` into kernel launch parameters and bindings.
-- **``LowerTIRxCleanup``** runs the ``LayoutApplier``: it resolves every
-  ``TileLayout``-typed buffer access into concrete physical address arithmetic
-  (``addr = data + elem_offset + layout.apply(*coord, shape=shape)``), replaces
-  layout-aware buffer parameters with physical views, and removes explicit
-  buffer offsets.
-
-After ``LowerTIRx``, tile primitives and ``TileLayout`` indirection are gone,
-and execution-scope IDs have been resolved. Some opaque TIRx constructs still
-remain; the later ``LowerTIRxOpaque`` pass converts those before
+- **``TilePrimitiveDispatch``** chooses concrete implementations for tile
+  operations. TIRx represents operations such as ``copy``, ``gemm``, and
+  ``reduction`` as ``TilePrimitiveCall`` nodes; this pass selects a backend
+  implementation for each one. It also turns abstract execution-scope
+  identifiers such as ``T.cta_id`` and ``T.thread_id`` into kernel-launch
+  parameters and thread bindings.
+- **``LowerTIRxCleanup``** maps logical coordinates to physical indices. It
+  applies supported logical layouts to buffer accesses so later passes can work
+  directly with concrete index expressions.
+
+After ``LowerTIRx``, tile operations have been replaced by their selected
+implementations, logical layouts have become physical indices, and abstract
+identifiers such as ``T.cta_id`` and ``T.thread_id`` have become thread
+bindings. Thread-binding loops and TIRx-specific loop annotations may still
+remain; ``LowerTIRxOpaque`` normalizes those structures before
 ``tirx.transform.FlattenBuffer`` flattens ordinary TIR buffer accesses.
 
-End-to-end IR evolution
------------------------
+Compiling a Simple Kernel to CUDA
+---------------------------------
+
+The following scale kernel illustrates two transformations: how ``T.cta_id``
+and ``T.thread_id`` become concrete thread identifiers, and how one TIRx
+function is split into a CPU-side launcher and a GPU kernel. The kernel processes
+1,024 elements using 4 CUDA thread blocks (CTAs), with 256 threads per CTA.
 
-Take a one-line scale kernel:
+**1. TIRx source uses abstract thread identifiers.**
 
 .. code-block:: python
 
+    import tvm
+    from tvm.script import tirx as T
+
     @T.prim_func
     def scale(A_ptr: T.handle, B_ptr: T.handle):
-        A = T.match_buffer(A_ptr, (256,), "float32")
-        B = T.match_buffer(B_ptr, (256,), "float32")
-        T.device_entry(); bx = T.cta_id([1]); tx = T.thread_id([256])
-        B[tx] = A[tx] * T.float32(2.0)
-
-This simple 1-D kernel has no nontrivial ``TileLayout``; it chiefly shows how
-``LowerTIRx`` turns scope IDs into real thread axes. The core body looks like the
-following excerpt. Buffer declarations and an unused warp-ID binding are omitted;
-``A_1`` and ``B_1`` are the materialized physical views:
+        A = T.match_buffer(A_ptr, (1024,), "float32")
+        B = T.match_buffer(B_ptr, (1024,), "float32")
+        T.device_entry()
+        bx = T.cta_id([4])
+        tx = T.thread_id([256])
+        B[bx * 256 + tx] = A[bx * 256 + tx] * T.float32(2.0)
+
+``T.device_entry()`` marks the entry into GPU code. ``LowerTIRx`` uses the
+marker to establish the corresponding thread bindings; the later
+``SplitHostDevice`` pass extracts the resulting device region into a separate
+kernel. ``T.cta_id([4])`` specifies 4 CTAs along x, while
+``T.thread_id([256])`` specifies 256 threads per CTA. At this point, ``bx`` and
+``tx`` are still abstract TIRx identifiers.
+
+**2. ``LowerTIRx`` lowers the abstract identifiers to TIR thread bindings.** It
+binds ``bx`` to ``blockIdx.x`` and ``tx`` to ``threadIdx.x``. Omitting buffer
+declarations, the core computation is equivalent to:
 
 .. code-block:: python
 
-    # match_buffer / decl_buffer declarations omitted
-    with T.launch_thread("blockIdx.x", 1) as blockIdx_x:
-        threadIdx_x = T.launch_thread("threadIdx.x", 256)
-        bx: T.let = blockIdx_x
-        tx: T.let = threadIdx_x
-        B_1[threadIdx_x] = A_1[threadIdx_x] * T.float32(2.0)
+    with T.launch_thread("blockIdx.x", 4) as bx:
+        tx = T.launch_thread("threadIdx.x", 256)
+        B[bx * 256 + tx] = A[bx * 256 + tx] * T.float32(2.0)
+
+This is still TIR, not CUDA source code. The excerpt retains only the important
+mapping; the next section shows how to print the complete compiler output.
 
-``SplitHostDevice`` then turns the single function into a host launcher and a
-device kernel. ``MakePackedAPI`` later rewrites the host launcher to TVM's
-packed-function ABI:
+**3. Later passes split host/device code and generate CUDA.** The compiler starts
+with one TIRx function. After ``LowerTIRx`` establishes thread bindings and a
+device region, ``SplitHostDevice`` produces two TIR functions (PrimFuncs):
 
 .. code-block:: text
 
-    @I.ir_module
-    class Module:
-        def main(...):          # host: packed-API launcher (computes the grid/block, launches)
-            ...
-        def scale_kernel(...):  # device: the __global__ body, run on the GPU
+    host launcher (generated from scale)
+      `-- launch scale_kernel with gridDim.x = 4 and blockDim.x = 256
+
+    device scale_kernel
+      `-- each GPU thread multiplies one input element by 2
 
-The CUDA backend then renders ``scale_kernel`` to the ``__global__`` function
-(``B_ptr[threadIdx.x] = A_ptr[threadIdx.x] * 2.0f``).
+The host function retains the kernel-launch logic, while the device function
+retains the elementwise computation. ``MakePackedAPI`` then adapts the host
+function to the uniform calling convention used by the TVM runtime. The device
+function proceeds to the CUDA backend, which generates code equivalent to:
+
+.. code-block:: cuda
+
+    __global__ void scale_kernel(float* A, float* B) {
+        int i = blockIdx.x * 256 + threadIdx.x;
+        B[i] = A[i] * 2.0f;
+    }
+
+In short, TIRx describes the thread organization and computation,
+``LowerTIRx`` turns abstract identifiers into TIR thread bindings,
+``SplitHostDevice`` separates CPU-side launch logic from GPU-side computation,
+and the CUDA backend finally emits CUDA source code.
+
+No bounds check is needed here because ``4 * 256`` is exactly 1,024. For a
+general length ``N``, choose the CTA count with ceiling division and guard the
+kernel body with ``i < N``.
 
 Inspecting intermediate IR and generated code
 ----------------------------------------------
 
-You can run any prefix of the pipeline by hand to inspect a stage — this is how the
-IR snippets across these docs were produced:
+To inspect an intermediate IR, run only the first few passes and stop before the
+rest of the pipeline. The following code first places ``scale`` in an
+``IRModule`` under the global name ``main``. The CUDA target selects the GPU
+backend, while ``with_host("llvm")`` selects LLVM for the CPU-side launcher.
+``BindTarget`` attaches both choices to the module, after which we run only
+``LowerTIRx``:
 
 .. code-block:: python
 
     from tvm.tirx import transform as TT
 
-    target = tvm.target.Target("cuda")
-    mod = TT.BindTarget(target.with_host("llvm"))(tvm.IRModule({"main": scale}))
-    mod = TT.LowerTIRx()(mod)         # tile primitives dispatched, layouts applied
-    print(mod.script())               # inspect the lowered TIRx IR
+    target = tvm.target.Target("cuda").with_host("llvm")
+    mod = tvm.IRModule({"main": scale})
+    mod = TT.BindTarget(target)(mod)
+    mod = TT.LowerTIRx()(mod)         # run LowerTIRx to lower abstract thread IDs
+    print(mod.script())               # inspect the IR after LowerTIRx
+
+The output should contain thread bindings for ``blockIdx.x`` and ``threadIdx.x``;
+the original ``T.cta_id`` and ``T.thread_id`` calls should be gone.
 
-Or compile the whole module and read the generated CUDA:
+To inspect the final CUDA, run the complete pipeline. The host module in this
+example imports exactly one device module, so ``imports[0]`` is the generated
+CUDA module, and ``inspect_source()`` returns its source code:
 
 .. code-block:: python
 
     exe = tvm.compile(tvm.IRModule({"main": scale}), target=target, tir_pipeline="tirx")
-    print(exe.mod.imports[0].inspect_source())
+    cuda_mod = exe.mod.imports[0]
+    print(cuda_mod.inspect_source())
+
+The generated code should contain ``blockIdx.x``, ``threadIdx.x``, and the
+elementwise multiplication that doubles each input value.
diff --git a/zh/tirx_guide/arch/index.rst b/zh/tirx_guide/arch/index.rst
index 0c0ed23b..0dabac96 100644
--- a/zh/tirx_guide/arch/index.rst
+++ b/zh/tirx_guide/arch/index.rst
@@ -20,4 +20,4 @@
 编译器内部机制
 ==============
 
-完整内容已移至 :ref:`TIRx 编译器内部机制：编译与 Lowering 流水线 <chap_arch>`。
+完整内容已移至 :ref:`chap_arch`。
diff --git a/zh/tirx_guide/arch/lowering_pipeline.rst b/zh/tirx_guide/arch/lowering_pipeline.rst
index af60d211..f06c4b00 100644
--- a/zh/tirx_guide/arch/lowering_pipeline.rst
+++ b/zh/tirx_guide/arch/lowering_pipeline.rst
@@ -17,50 +17,48 @@
 
 .. _chap_arch:
 
-TIRx 编译器内部机制：编译与 Lowering 流水线
-================================================
+TIRx 编译器内部机制
+===================
 
-``tvm.compile(mod, target, tir_pipeline="tirx")`` 会把编写好的 TIRx module
-转换成 host launcher 和 device code。这个过程并非一次完成：编译器先处理
-TIRx 特有的结构，再用通用的 TIRx 规范化与合法化 passes 处理结果，最后拆分
-module 并交给后端生成代码。
+``tvm.compile(mod, target, tir_pipeline="tirx")`` 接收一个 TIRx module，最终生成两部分代码：CPU 端的启动函数负责准备参数并启动 GPU，GPU 端的 kernel 负责执行计算。编译器不是一步完成这项转换，而是依次运行多个编译步骤。每个步骤称为一个 pass，负责对 IR 做一类特定的转换、检查或标注。
 
-完整顺序定义在 `compilation_pipeline.py
+TIRx 的完整 pass 顺序定义在 Apache TVM 源码中的 `compilation_pipeline.py
 <https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/compilation_pipeline.py>`_。
-本页先说明这条流水线在 ``tvm.compile`` 中的位置，再按阶段解释各个 pass
-改变了什么，以及 host 与 device 两条路径从哪里分开。
 
 整体编译路径
 ------------
 
-``tvm.compile`` 首先绑定 target，再运行下面的 module-level **tirx
-pipeline**。随后，host 和 device functions 分别经过 finalization passes，
-device function 最终交给 CUDA code generator：
+``target`` 指定代码要在哪种硬件和后端上运行。下面的例子在 GPU 端使用 CUDA，在 CPU 端使用 LLVM：``tvm.compile`` 先将 target 信息写入 module，再运行模块级的 **tirx pipeline**。``tirx_pipeline`` 拆出 CPU 端的 host 函数和 GPU 端的 device 函数后，两者分别经过最后一轮面向具体 target 的转换，再交给相应的代码生成器：
 
 .. code-block:: text
 
-    authored TIRx  ──BindTarget──▶  tirx_pipeline  ──▶  host func  ──host finalize──▶  C/LLVM
-                                          │
-                                          └──────────▶  device func ──device finalize──▶  CUDA
+    编写好的 TIRx
+          │ BindTarget
+          ▼
+    tirx_pipeline
+    （SplitHostDevice 在其中拆分两条路径）
+          ├── host PrimFunc   ──host finalization──▶ C/LLVM
+          └── device PrimFunc ─device finalization─▶ CUDA
 
-``tirx_pipeline`` 的 Pass 执行顺序
------------------------------------
+``PrimFunc`` 是 TIR 中的函数表示。上图中的 host PrimFunc 就是 CPU 端的启动函数，device PrimFunc 则是 GPU 上执行的 kernel；``finalization`` 表示代码生成之前针对具体 target 所做的最后几步转换。
 
-``tirx_pipeline`` 按下表中的 19 个步骤组织。公共子表达式消除是可选项，
-vectorization 和 unrolling 的行为也可以通过 ``PassContext`` 控制：
+``tirx_pipeline`` 的 pass 顺序
+-------------------------------
+
+下表按照实际执行顺序列出 ``tirx_pipeline`` 中的 19 个步骤。ABI 是函数之间的调用约定；表中的 ABI passes 负责把普通 TIR 函数改造成 runtime 能够调用的形式。``PassContext`` 是控制编译选项的配置对象：公共子表达式消除可以关闭，向量化和循环展开的行为也可以通过它调整。
 
 .. list-table::
    :header-rows: 1
    :widths: 6 24 24 46
 
    * - #
-     - 阶段
+     - 类别
      - Pass
      - 作用
    * - 1
      - TIRx lowering
      - ``LowerTIRx``
-     - 完成 TIRx 的核心转换，详见下方 `LowerTIRx 内部做了什么`_
+     - 完成 TIRx 的核心转换，详见下方 `LowerTIRx 的内部组成`_
    * - 2
      - TIR 规范化
      - ``UnifyThreadBinding``
@@ -69,12 +67,11 @@ vectorization 和 unrolling 的行为也可以通过 ``PassContext`` 控制：
    * - 3
      - TIR 规范化
      - ``StmtSimplify``
-     - 使用 arithmetic analyzer 简化 statement 中的算术表达式
+     - 简化 IR 中的算术表达式
    * - 4
      - TIR 规范化
      - ``LowerTIRxOpaque``
-     - 处理剩余的 opaque constructs，包括 thread-binding loops、unit loops
-       和 pragma annotations
+     - 转换绑定到线程轴的循环，消除未标注的长度为 1 的循环，并规范化循环 pragma
    * - 5
      - TIR 规范化
      - ``FlattenBuffer``
@@ -82,25 +79,24 @@ vectorization 和 unrolling 的行为也可以通过 ``PassContext`` 控制：
    * - 6
      - 计算合法化
      - ``BF16ComputeLegalize``
-     - 将 ``bfloat16`` 计算改写为合法形式，其中计算会提升到 f32
+     - target 不原生支持 ``bfloat16`` 计算时，将其提升到 ``float32`` 并改写为合法形式
    * - 7
      - TIR 规范化
      - ``NarrowDataType(32)``
-     - 在能够证明安全时，将 index 和 loop 的 scalar ``Expr`` type 缩窄为 32 bits
+     - 在能够证明安全时，将索引表达式和循环变量缩窄至 32 位
    * - 8
-     - Loop lowering
+     - 循环转换
      - ``VectorizeLoop``
-     - 将 ``T.vectorized`` loops 转换为 vector operations；设置
-       ``tir.disable_vectorize`` 时，改为将这些 loops scalarize
+     - 将 ``T.vectorized`` 循环转换为向量操作；设置 ``tir.disable_vectorize`` 时，则改写为普通标量循环
    * - 9
-     - Loop lowering
+     - 循环转换
      - ``UnrollLoop``
-     - 展开标记为 ``T.unroll`` 的 loops；普通常量 loops 只有在相应配置或
+     - 展开标记为 ``T.unroll`` 的循环；普通常量循环只有在相应配置或
        pragma 启用时才会自动展开
    * - 10
      - TIR 规范化
      - ``StmtSimplify``
-     - Vectorize 和 unroll 暴露出更多常量后，再次执行简化
+     - 向量化和循环展开后会出现更多可简化的常量，再次执行简化
    * - 11
      - TIR 规范化
      - ``CommonSubexprElim``
@@ -108,54 +104,52 @@ vectorization 和 unrolling 的行为也可以通过 ``PassContext`` 控制：
    * - 12
      - 计算合法化
      - ``FP8ComputeLegalize``
-     - 将 ``float8`` 计算改写为合法形式
+     - target 不原生支持 ``float8`` 计算时，将其提升为受支持的类型
+       （默认为 ``float32``）
    * - 13
      - 校验与 ABI
      - ``VerifyMemory``
-     - 检查 host 代码没有直接解引用 device memory
+     - 确保 host 代码不会直接解引用 device memory
    * - 14
      - 校验与 ABI
      - ``AnnotateEntryFunc``
-     - 将唯一 function 标记为入口；对于多 function module，则标记其中唯一
-       对外可见的 PrimFunc
+     - 只有一个 PrimFunc 时直接将其标记为入口；有多个 PrimFunc 时，则标记其中唯一对外可见的函数
    * - 15
      - 校验与 ABI
      - ``SplitHostDevice``
-     - 识别 device regions，拆分 host 与 device PrimFuncs，并将 host 侧调用
-       转换为 kernel-launch ABI
+     - 识别 device regions，拆分 host 与 device PrimFuncs，并将 host 侧调用转换为 kernel-launch ABI
    * - 16
      - 校验与 ABI
      - ``LowerIket``
-     - 普通 build 中移除 frontend-only NVIDIA IKET annotations；IRModule 显式
-       启用 IKET 时则生成 IKET metadata 和 placeholders
+     - 普通 build 中移除 NVIDIA IKET annotations；启用 IKET 时则将其转换为 tracing 所需的形式
    * - 17
      - 校验与 ABI
      - ``MakePackedAPI``
-     - 将 host function 改写为 TVM launcher 使用的 packed-function ABI
+     - 将 host function 改写为 TVM runtime 通过 packed-function ABI 调用的形式
    * - 18
-     - Storage 合法化
+     - 存储合法化
      - ``FP8StorageLegalize``
-     - 将 ``float8`` storage 转换为 ``uint8`` container
+     - target 不原生支持 ``float8`` storage 时，改用 ``uint8`` container 保存
    * - 19
-     - Storage 合法化
+     - 存储合法化
      - ``BF16StorageLegalize``
-     - 将 ``bfloat16`` storage 转换为 ``uint16`` container
+     - target 不原生支持 ``bfloat16`` storage 时，改用 ``uint16`` container 保存
 
-Host 与 Device Finalization
----------------------------
+Host 与 Device 的后续处理
+-------------------------
 
-上面列出的 19 个步骤组成 ``tirx_pipeline``。这条 module-level pipeline
-结束后，``tvm.compile`` 会根据 function 类型分别执行 finalization：
+上面列出的 19 个步骤组成 ``tirx_pipeline``。这条模块级 pipeline
+结束后，``tvm.compile`` 会根据函数类型分别执行 finalization：
 
 - **host**：``LowerTVMBuiltin`` 处理 ``tvm_*`` builtins，``LowerIntrin``
-  处理 target-specific intrinsics。
+  处理面向具体 target 的 intrinsics。
 - **device**：``LowerWarpMemory`` 将 warp-scoped buffers 转换为
   shuffles，随后执行 ``StmtSimplify`` 和 ``LowerIntrin``。
 
-``LowerTIRx`` 内部做了什么
----------------------------
+``LowerTIRx`` 的内部组成
+------------------------
 
-正常编译时，``LowerTIRx`` 本身由两个 passes 组成，定义在
+``LowerTIRx`` 主要完成两个任务：为 tile-level 操作选择具体实现，以及把逻辑数据布局转换成实际的内存索引。它的核心转换由下面两个 passes 组成，定义在 Apache TVM 源码中的
 `lower_tirx.cc
 <https://github.com/apache/tvm/blob/v0.26.0/src/tirx/transform/lower_tirx.cc>`_：
 
@@ -163,80 +157,108 @@ Host 与 Device Finalization
 
     LowerTIRx = Sequential([ TilePrimitiveDispatch, LowerTIRxCleanup ])
 
-- **``TilePrimitiveDispatch``** 根据 backend dispatch，为每个
-  ``TilePrimitiveCall``（``copy``、``gemm``、``reduction`` 等）选择具体实现，
-  同时把 ``T.cta_id``、``T.thread_id`` 等 execution-scope IDs 解析为 kernel
-  launch parameters 和对应的 bindings。
-- **``LowerTIRxCleanup``** 运行 ``LayoutApplier``，将使用
-  ``TileLayout`` 的 buffer access 变成具体的物理地址计算
-  （``addr = data + elem_offset + layout.apply(*coord, shape=shape)``），把带 layout
-  的 buffer parameters 替换成物理 views，并移除显式的 buffer offsets。
+- **``TilePrimitiveDispatch``** 为 tile 操作选择具体实现。TIRx 中的 ``copy``、
+  ``gemm``、``reduction`` 等操作以 ``TilePrimitiveCall`` 表示；这个 pass 根据
+  backend 选择对应实现。它还会把 ``T.cta_id``、``T.thread_id`` 等抽象的执行范围编号转换成 kernel launch 参数和线程绑定。
+- **``LowerTIRxCleanup``** 将逻辑坐标转换为物理索引。它把支持的逻辑 layout 应用到 buffer access 上，使后续 passes 可以直接处理具体的索引表达式。
+
+完成 ``LowerTIRx`` 后，tile 操作已经换成选定的底层实现，逻辑 layout 也已经落实为物理索引，``T.cta_id`` 和 ``T.thread_id`` 等抽象编号则变成了线程绑定。此时仍可能保留 thread-binding loops 和 TIRx 特有的 loop annotations；后续的
+``LowerTIRxOpaque`` 会规范化这些结构，再由 ``tirx.transform.FlattenBuffer``
+展平普通 TIR 中的 buffer access。
 
-完成 ``LowerTIRx`` 后，tile primitives 和 ``TileLayout`` 间接层已经消失，
-execution-scope IDs 也已经解析。此时仍有少量 opaque TIRx constructs；后续的
-``LowerTIRxOpaque`` 会先处理这些结构，后续的
-``tirx.transform.FlattenBuffer`` pass 再展平普通 TIR 中的 buffer access。
+一个简单 Kernel 的编译过程
+--------------------------
 
-端到端 IR 演化示例
-------------------
+下面用一个 scale kernel 观察两件事：``T.cta_id`` 和 ``T.thread_id`` 怎样落实为具体的线程编号，以及一个 TIRx 函数如何拆成 CPU 端的启动函数与 GPU 上执行的 kernel。这个 kernel 处理 1,024 个元素，使用 4 个 CUDA thread blocks（CTA），每个 CTA 包含 256 个线程。
 
-以下面的 scale kernel 为例：
+**1. TIRx 源码使用抽象的线程编号。**
 
 .. code-block:: python
 
+    import tvm
+    from tvm.script import tirx as T
+
     @T.prim_func
     def scale(A_ptr: T.handle, B_ptr: T.handle):
-        A = T.match_buffer(A_ptr, (256,), "float32")
-        B = T.match_buffer(B_ptr, (256,), "float32")
-        T.device_entry(); bx = T.cta_id([1]); tx = T.thread_id([256])
-        B[tx] = A[tx] * T.float32(2.0)
+        A = T.match_buffer(A_ptr, (1024,), "float32")
+        B = T.match_buffer(B_ptr, (1024,), "float32")
+        T.device_entry()
+        bx = T.cta_id([4])
+        tx = T.thread_id([256])
+        B[bx * 256 + tx] = A[bx * 256 + tx] * T.float32(2.0)
 
-这个简单的一维 kernel 没有非平凡的 ``TileLayout``，主要用来展示
-``LowerTIRx`` 如何将 scope IDs 转换成真实的 thread axes。下面只摘录核心
-body，省略 buffer declarations 和未使用的 warp-ID binding；``A_1`` 与
-``B_1`` 是生成的物理 views：
+``T.device_entry()`` 标记 GPU 代码的入口。``LowerTIRx`` 根据这个标记建立线程绑定；后面的 ``SplitHostDevice`` 再从所得设备代码区域（device region）中拆出单独的 device kernel。``T.cta_id([4])`` 表示 x 方向有 4 个 CTA，
+``T.thread_id([256])`` 表示每个 CTA 有 256 个线程。这里的 ``bx`` 和 ``tx``
+仍是 TIRx 提供的抽象编号。
+
+**2. ``LowerTIRx`` 将抽象编号转换为 TIR 线程绑定。** 它把 ``bx`` 和 ``tx``
+分别绑定到 ``blockIdx.x`` 与 ``threadIdx.x``。省略 buffer declarations 后，核心计算等价于：
 
 .. code-block:: python
 
-    # 省略 match_buffer / decl_buffer declarations
-    with T.launch_thread("blockIdx.x", 1) as blockIdx_x:
-        threadIdx_x = T.launch_thread("threadIdx.x", 256)
-        bx: T.let = blockIdx_x
-        tx: T.let = threadIdx_x
-        B_1[threadIdx_x] = A_1[threadIdx_x] * T.float32(2.0)
+    with T.launch_thread("blockIdx.x", 4) as bx:
+        tx = T.launch_thread("threadIdx.x", 256)
+        B[bx * 256 + tx] = A[bx * 256 + tx] * T.float32(2.0)
+
+这里仍然是 TIR，还不是 CUDA 源码。这段代码只保留了关键映射，不是编译器输出的完整 IR；下一节会给出打印完整结果的命令。
 
-``SplitHostDevice`` 随后将单个 function 拆成 host launcher 和 device kernel，
-``MakePackedAPI`` 再将 host launcher 转换为 TVM 的 packed-function ABI：
+**3. 后续 passes 拆分 host/device，并生成 CUDA。** 编译开始时只有一个 TIRx
+函数。``LowerTIRx`` 生成线程绑定和 device region 后，``SplitHostDevice`` 将其拆成两个 TIR 函数（PrimFunc）：
 
 .. code-block:: text
 
-    @I.ir_module
-    class Module:
-        def main(...):          # host: packed-API launcher (computes the grid/block, launches)
-            ...
-        def scale_kernel(...):  # device: the __global__ body, run on the GPU
+    host launcher（由 scale 生成）
+      └── 启动 scale_kernel，gridDim.x = 4，blockDim.x = 256
+
+    device scale_kernel
+      └── 每个 GPU 线程将一个输入元素乘以 2
 
-CUDA backend 随后将 ``scale_kernel`` 生成 ``__global__`` function：
-``B_ptr[threadIdx.x] = A_ptr[threadIdx.x] * 2.0f``。
+host 函数保存 kernel 的启动逻辑；device 函数保存真正的逐元素计算。
+``MakePackedAPI`` 随后将 host 函数转换为 TVM runtime 使用的统一调用形式。Device 函数则交给 CUDA backend，生成与下面代码等价的 CUDA
+kernel：
+
+.. code-block:: cuda
+
+    __global__ void scale_kernel(float* A, float* B) {
+        int i = blockIdx.x * 256 + threadIdx.x;
+        B[i] = A[i] * 2.0f;
+    }
+
+整个过程可以概括为：TIRx 描述线程组织和计算，``LowerTIRx`` 将抽象编号落实为
+TIR 线程绑定，``SplitHostDevice`` 分开 CPU 端的启动逻辑与 GPU 端的计算，CUDA
+backend 最后生成 CUDA 源码。
+
+这里不需要边界判断，因为 ``4 * 256`` 恰好等于 1,024。处理一般长度 ``N``
+时，需要向上取整得到 CTA 数量，并在 kernel 中判断 ``i < N``。
 
 检查中间 IR 与生成代码
 ----------------------
 
-可以手动运行 pipeline 的任意前缀，检查某个阶段的 IR。本书中的 IR 片段也是
-用这种方式生成的：
+为了查看中间 IR，可以只运行完整 pipeline 最前面的几步，然后停下来打印结果。下面先把 ``scale`` 以全局名 ``main`` 放入 ``IRModule``。CUDA target 指定 GPU
+端生成 CUDA，``with_host("llvm")`` 则指定 CPU 端生成 LLVM 代码。
+``BindTarget`` 将这组 target 信息写入 module，随后只运行 ``LowerTIRx``：
 
 .. code-block:: python
 
     from tvm.tirx import transform as TT
 
-    target = tvm.target.Target("cuda")
-    mod = TT.BindTarget(target.with_host("llvm"))(tvm.IRModule({"main": scale}))
-    mod = TT.LowerTIRx()(mod)         # tile primitives dispatched, layouts applied
-    print(mod.script())               # inspect the lowered TIRx IR
+    target = tvm.target.Target("cuda").with_host("llvm")
+    mod = tvm.IRModule({"main": scale})
+    mod = TT.BindTarget(target)(mod)
+    mod = TT.LowerTIRx()(mod)         # 运行 LowerTIRx，转换抽象线程编号
+    print(mod.script())               # 查看 LowerTIRx 之后的 IR
+
+输出中应该能看到 ``blockIdx.x`` 和 ``threadIdx.x`` 对应的线程绑定，而原来的
+``T.cta_id`` 与 ``T.thread_id`` 已经消失。
 
-也可以编译完整 module，再查看生成的 CUDA：
+要查看最终 CUDA，可以运行完整 pipeline。这里的 host module 只导入了一个
+device module，因此 ``imports[0]`` 就是生成的 CUDA module；
+``inspect_source()`` 返回它的源码：
 
 .. code-block:: python
 
     exe = tvm.compile(tvm.IRModule({"main": scale}), target=target, tir_pipeline="tirx")
-    print(exe.mod.imports[0].inspect_source())
+    cuda_mod = exe.mod.imports[0]
+    print(cuda_mod.inspect_source())
+
+生成的代码中应该能找到 ``blockIdx.x``、``threadIdx.x``，以及将每个输入元素乘以 2 的计算。

From c4f6611e20ce808783d22444d8b0b301cef244ec Mon Sep 17 00:00:00 2001
From: tlopex <820958424@qq.com>
Date: Sat, 15 Aug 2026 03:12:20 -0400
Subject: [PATCH 3/4] Restore TIRx compiler internals chapter hierarchy

---
 index.md                                 |  2 +-
 tirx_guide/arch/index.rst                | 16 ++++++++++++----
 tirx_guide/arch/lowering_pipeline.rst    |  6 ++----
 zh/index.md                              |  2 +-
 zh/tirx_guide/arch/index.rst             | 13 +++++++++----
 zh/tirx_guide/arch/lowering_pipeline.rst |  6 ++----
 6 files changed, 27 insertions(+), 18 deletions(-)

diff --git a/index.md b/index.md
index 5e903ef6..ff398443 100644
--- a/index.md
+++ b/index.md
@@ -87,6 +87,6 @@ chapter_flash_attention/index
 
 appendix/index
 tirx_guide/language_reference/index
-TIRx Compiler Internals <tirx_guide/arch/lowering_pipeline>
+tirx_guide/arch/index
 appendix/debugging_warp_specialized
 ```
diff --git a/tirx_guide/arch/index.rst b/tirx_guide/arch/index.rst
index 5a213844..98a2b6b5 100644
--- a/tirx_guide/arch/index.rst
+++ b/tirx_guide/arch/index.rst
@@ -15,9 +15,17 @@
     specific language governing permissions and limitations
     under the License.
 
-:orphan:
+.. _chap_arch:
 
-Compiler Internals
-==================
+TIRx Compiler Internals
+=======================
 
-The compiler-internals guide is now available at :ref:`chap_arch`.
+This section explains how the TIRx compiler lowers an authored module into a
+CPU-side launcher and GPU device code. It follows the compilation pipeline from
+high-level TIRx constructs through host/device separation and CUDA code
+generation.
+
+.. toctree::
+   :maxdepth: 1
+
+   lowering_pipeline
diff --git a/tirx_guide/arch/lowering_pipeline.rst b/tirx_guide/arch/lowering_pipeline.rst
index 792fdcbd..49efa927 100644
--- a/tirx_guide/arch/lowering_pipeline.rst
+++ b/tirx_guide/arch/lowering_pipeline.rst
@@ -15,10 +15,8 @@
     specific language governing permissions and limitations
     under the License.
 
-.. _chap_arch:
-
-TIRx Compiler Internals
-=======================
+TIRx Lowering Pipeline
+======================
 
 ``tvm.compile(mod, target, tir_pipeline="tirx")`` takes a TIRx module and
 eventually produces two pieces of code: a CPU-side launcher that prepares the
diff --git a/zh/index.md b/zh/index.md
index b0f43df4..b2567331 100644
--- a/zh/index.md
+++ b/zh/index.md
@@ -64,6 +64,6 @@ chapter_flash_attention/index
 
 appendix/index
 tirx_guide/language_reference/index
-TIRx 编译器内部机制 <tirx_guide/arch/lowering_pipeline>
+tirx_guide/arch/index
 appendix/debugging_warp_specialized
 ```
diff --git a/zh/tirx_guide/arch/index.rst b/zh/tirx_guide/arch/index.rst
index 0dabac96..498d0ae0 100644
--- a/zh/tirx_guide/arch/index.rst
+++ b/zh/tirx_guide/arch/index.rst
@@ -15,9 +15,14 @@
     specific language governing permissions and limitations
     under the License.
 
-:orphan:
+.. _chap_arch:
 
-编译器内部机制
-==============
+TIRx 编译器内部机制
+===================
 
-完整内容已移至 :ref:`chap_arch`。
+本节介绍 TIRx 编译器如何将编写好的 module 转换成 CPU 端的启动函数和 GPU 端的 device code，并沿着编译流水线说明 TIRx 高层结构、host/device 拆分以及 CUDA 代码生成之间的关系。
+
+.. toctree::
+   :maxdepth: 1
+
+   lowering_pipeline
diff --git a/zh/tirx_guide/arch/lowering_pipeline.rst b/zh/tirx_guide/arch/lowering_pipeline.rst
index f06c4b00..672ef298 100644
--- a/zh/tirx_guide/arch/lowering_pipeline.rst
+++ b/zh/tirx_guide/arch/lowering_pipeline.rst
@@ -15,10 +15,8 @@
     specific language governing permissions and limitations
     under the License.
 
-.. _chap_arch:
-
-TIRx 编译器内部机制
-===================
+TIRx 编译流水线
+===============
 
 ``tvm.compile(mod, target, tir_pipeline="tirx")`` 接收一个 TIRx module，最终生成两部分代码：CPU 端的启动函数负责准备参数并启动 GPU，GPU 端的 kernel 负责执行计算。编译器不是一步完成这项转换，而是依次运行多个编译步骤。每个步骤称为一个 pass，负责对 IR 做一类特定的转换、检查或标注。
 

From ff854d4b0e1d749d363223c0102f5901a50aa326 Mon Sep 17 00:00:00 2001
From: tlopex <820958424@qq.com>
Date: Sat, 15 Aug 2026 03:22:21 -0400
Subject: [PATCH 4/4] Polish compiler internals titles and source links

---
 tirx_guide/arch/index.rst                | 4 ++--
 tirx_guide/arch/lowering_pipeline.rst    | 4 ++--
 zh/tirx_guide/arch/index.rst             | 4 ++--
 zh/tirx_guide/arch/lowering_pipeline.rst | 4 ++--
 4 files changed, 8 insertions(+), 8 deletions(-)

diff --git a/tirx_guide/arch/index.rst b/tirx_guide/arch/index.rst
index 98a2b6b5..0409bd47 100644
--- a/tirx_guide/arch/index.rst
+++ b/tirx_guide/arch/index.rst
@@ -17,8 +17,8 @@
 
 .. _chap_arch:
 
-TIRx Compiler Internals
-=======================
+Compiler Internals
+==================
 
 This section explains how the TIRx compiler lowers an authored module into a
 CPU-side launcher and GPU device code. It follows the compilation pipeline from
diff --git a/tirx_guide/arch/lowering_pipeline.rst b/tirx_guide/arch/lowering_pipeline.rst
index 49efa927..9e71142e 100644
--- a/tirx_guide/arch/lowering_pipeline.rst
+++ b/tirx_guide/arch/lowering_pipeline.rst
@@ -25,7 +25,7 @@ computation. The compiler reaches that result through an ordered series of
 passes. Each pass performs a particular transformation, validation, or
 annotation on the IR.
 
-The exact pass order is defined in Apache TVM's `compilation_pipeline.py
+The exact pass order is defined in Apache TVM's `python/tvm/tirx/compilation_pipeline.py
 <https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/compilation_pipeline.py>`_.
 
 The overall compilation path
@@ -175,7 +175,7 @@ Inside ``LowerTIRx``
 ``LowerTIRx`` has two main jobs: choosing concrete implementations for tile-level
 operations, and turning logical data layouts into physical memory indices. Its
 core transformation is the following two-pass sequence, defined in Apache TVM's
-`lower_tirx.cc
+`src/tirx/transform/lower_tirx.cc
 <https://github.com/apache/tvm/blob/v0.26.0/src/tirx/transform/lower_tirx.cc>`_:
 
 .. code-block:: text
diff --git a/zh/tirx_guide/arch/index.rst b/zh/tirx_guide/arch/index.rst
index 498d0ae0..db9145de 100644
--- a/zh/tirx_guide/arch/index.rst
+++ b/zh/tirx_guide/arch/index.rst
@@ -17,8 +17,8 @@
 
 .. _chap_arch:
 
-TIRx 编译器内部机制
-===================
+编译器内部机制
+==============
 
 本节介绍 TIRx 编译器如何将编写好的 module 转换成 CPU 端的启动函数和 GPU 端的 device code，并沿着编译流水线说明 TIRx 高层结构、host/device 拆分以及 CUDA 代码生成之间的关系。
 
diff --git a/zh/tirx_guide/arch/lowering_pipeline.rst b/zh/tirx_guide/arch/lowering_pipeline.rst
index 672ef298..d5ecd121 100644
--- a/zh/tirx_guide/arch/lowering_pipeline.rst
+++ b/zh/tirx_guide/arch/lowering_pipeline.rst
@@ -20,7 +20,7 @@ TIRx 编译流水线
 
 ``tvm.compile(mod, target, tir_pipeline="tirx")`` 接收一个 TIRx module，最终生成两部分代码：CPU 端的启动函数负责准备参数并启动 GPU，GPU 端的 kernel 负责执行计算。编译器不是一步完成这项转换，而是依次运行多个编译步骤。每个步骤称为一个 pass，负责对 IR 做一类特定的转换、检查或标注。
 
-TIRx 的完整 pass 顺序定义在 Apache TVM 源码中的 `compilation_pipeline.py
+TIRx 的完整 pass 顺序定义在 Apache TVM 源码中的 `python/tvm/tirx/compilation_pipeline.py
 <https://github.com/apache/tvm/blob/v0.26.0/python/tvm/tirx/compilation_pipeline.py>`_。
 
 整体编译路径
@@ -148,7 +148,7 @@ Host 与 Device 的后续处理
 ------------------------
 
 ``LowerTIRx`` 主要完成两个任务：为 tile-level 操作选择具体实现，以及把逻辑数据布局转换成实际的内存索引。它的核心转换由下面两个 passes 组成，定义在 Apache TVM 源码中的
-`lower_tirx.cc
+`src/tirx/transform/lower_tirx.cc
 <https://github.com/apache/tvm/blob/v0.26.0/src/tirx/transform/lower_tirx.cc>`_：
 
 .. code-block:: text