• 文档
• 教程
• 人工智能

本文是 AISystem 开源课程《编译后端优化》系列中"算子手工优化"一节的完整技术指南。文章以算子计算与调度为起点，围绕 Roofline 模型教你如何定量判断算子的计算瓶颈与访存瓶颈，随后梳理循环优化、指令优化、存储优化三大策略的核心思想，并深入讲解 TVM 调度 API 与 Triton 分块编程这两种主流 DSL 开发范式及其底层编译原理。读完本文，你将掌握"先分析瓶颈、再选择策略、最后借助 DSL 高效落地"的完整算子优化方法论。

从计算与调度说起：为什么需要手工优化

在进入优化细节之前，需要先建立两个基本概念：计算（Compute）与调度（Schedule）。如本系列前文 02OPSCompute.md 所述，计算描述实现算法的具体逻辑而不关心代码实现，调度则是对计算进行优化和控制的过程——通过指定执行顺序、内存布局、并行策略来实现性能优化。同一个矩阵乘法既可以朴素地三重循环实现，也可以通过循环分块、向量化等手段大幅加速，二者的性能差异可达几十倍甚至几百倍。这种"一对多"的映射空间被称为调度空间，AI 编译器后端优化的本质，就是在调度空间中为特定硬件寻找最优实现。

算子手工优化正是这一思想的直接实践：先通过 Roofline 模型判断瓶颈方向，再针对算子的循环结构、指令特征和访存模式实施定制化优化。下面我们沿此路径逐步展开。

计算分析：用 Roofline 模型定位瓶颈

优化算子前，首先要回答"瓶颈在哪里"：当前程序是计算瓶颈（Compute-Bound）还是访存瓶颈（Memory-Bound）？Roofline（屋顶线）模型正是回答这一问题的标准工具。它通过同时考察平台算力与内存带宽，给出程序在特定硬件上的理论性能上限，帮助开发者识别瓶颈并制定对应的优化策略。

四大关键指标

在 Roofline 分析之前，先明确四个基础指标：

• 计算量（Floating Point Operations）：程序完成一次完整计算发生的浮点运算个数，即时间复杂度，单位为 Flops。例如卷积层的计算量为 $M^2 \times K^2 \times C_{in} \times C_{out}$，其中 M 为输入特征图大小，K 为卷积核大小，C 为通道数。
• 访存量（Bytes）：程序完成一次完整计算发生的内存交换总量，即空间复杂度。理想情况下，访存量等于模型参数与输出特征图的内存占用总和（输入的内存占用计入上一个算子的输出占用），单位为 Byte；在 float32 类型下需乘以 4。对于卷积层，内存占用为 $K^2 \times C_{in} \times C_{out} + M^2 \times C_{out}$。
• 计算强度（Computational Intensity）：计算量除以访存量，表示每字节内存交换用于进行多少次浮点计算。计算强度越大，内存使用效率越高。
• 理论性能（Attainable Performance）：模型在计算平台上所能达到的每秒浮点运算次数上限，Roofline 模型给出的就是计算该指标的方法。

Roofline 模型：算力屋顶与带宽房檐

Roofline 模型是一种用于评估和分析高性能计算平台性能的有效工具。它通过分析计算量和访存量，确定在特定算力和带宽条件下计算任务所能达到的理论性能上限。其核心价值在于揭示硬件资源的限制，帮助开发者和研究者理解在现有计算平台约束下应用能够实现的理论性能上界。

模型由两部分构成：算力决定"屋顶"的高度（绿色线段，平台每秒最多可执行的浮点运算次数），带宽决定"房檐"的斜率（红色线段，带宽越大房檐越陡，内存受限区域的性能上限越高）。

据此可划分两种瓶颈状态：

• Compute-Bound：当算子的计算强度大于计算平台的计算强度上限时，算子处于计算瓶颈。此时算子已经利用了计算平台的全部算力，进一步优化方向是提升计算密度（如向量化、张量化）。
• Memory-Bound：当算子的计算强度小于计算平台的强度上限时，算子性能完全由平台带宽上限和模型自身计算强度决定。此时带宽越大（房檐越陡）、或模型计算强度越大，理论性能可呈线性增长，优化方向是提升数据局部性与访存效率。

实例：VGG16 与 MobileNet 在 1080Ti 上的对比

以（3,224,224）输入为例，Roofline 模型可以直观地区分两类典型网络：

• VGG16：前向传播计算量为 15 GFLOPs，访存量约 600 MB，计算强度为 25 FLOPs/Byte。
• MobileNet：计算量为 0.5 GFLOPs，访存量为 74 MB，计算强度仅 7 FLOPs/Byte。
• 平台参数（1080Ti GPU）：算力 11.3 TFLOP/s，带宽 484 GB/s，最大计算强度约 24。

从图中可以看出：MobileNet 处于Memory-Bound区域，在 1080Ti 上的理论性能只有 3.3 TFLOPs；VGG 处于Compute-Bound区域，可以完全利用 1080Ti 的全部算力。通过 Roofline 模型可以清晰地看到：当计算量和访存量增加时，性能提升会受到硬件算力和带宽的双重限制。这种分析对优化计算密集型和内存带宽密集型应用至关重要，可帮助开发者识别性能瓶颈并制定相应策略。

此外，Roofline 模型还可用于指导硬件设计和软件算法选择：若应用受限于内存带宽，开发者可考虑更高效的数据结构或算法以减少内存访问；硬件设计师也可利用它评估不同硬件配置对特定应用性能的影响。

需要强调的是，Roofline 模型给出的是理论性能上界。实际运行中，除算力和带宽外，缓存大小与速度、指令流水线、并行调度等因素都会影响最终性能。

优化策略：循环、指令、存储三大核心

明确瓶颈方向后，手工优化主要聚焦三大核心领域：循环优化、指令优化和存储优化。这些策略针对算子的计算特性与硬件资源特点量身定制，下文先做整体概述，其细节在本系列后续章节（04LoopOpt.md、05OtherOpt.md）中有更详尽的分析。

循环优化：利用规则化嵌套循环结构

AI 算子普遍具有高度规则化的多层嵌套循环结构，这为优化提供了丰富手段。以逐元素操作为例，如 ReLU、加法（Add）、乘法（Mul），可以在所有循环轴上迭代执行计算；即使是卷积（Conv）这类复杂算子，也可以通过七层嵌套循环实现。但直观的原生写法往往效率低下。

通过对算子的数据布局和内存访问特性进行分析，再相应调整循环结构，可以显著减少不必要的开销、更高效地利用硬件资源，从而降低延迟。常见的循环优化技术包括：

• 循环分块（Loop Blocking）：将大数据集切成可装入 Cache 的小块，提升数据复用、降低 Cache miss。
• 循环展开（Loop Unrolling）：将多次迭代展开，减少分支跳转与循环开销，提升指令级并行。
• 循环重排（Loop Reordering）：调整循环嵌套顺序，改善访存模式与数据局部性。
• 循环融合（Loop Fusion）：合并多个相邻循环，减少内存访问与循环开销。
• 循环拆分（Loop Splitting）：将一个循环拆成多个，分离控制流，创造更多优化机会。

仓库 04Backend/code 目录下提供了四组循环优化示例（伪代码形式），与本节策略一一对应：tiling 演示将A[i] += B[j]的 j 轴循环按 T 分块并把外层块循环上提；reorder 演示交换 i/j 两层循环顺序；fusion 演示将两个依赖相邻的循环融合为一个循环并处理首尾边界；spliting 演示把含条件分支的循环拆分为两个独立循环。这些变换通过精心设计，既能提升计算密度，又能改善数据局部性，最终实现性能提升。

指令优化：向量化与专用加速指令

现代处理器（如 Intel AVX-512、ARM SVE）提供强大的向量处理能力，允许单条指令同时操作多个数据点，从而减少指令数量和执行周期。针对特定硬件定制的指令集（如 NVIDIA Tensor Cores）则进一步增强并行处理能力，为深度学习张量运算提供专门优化。

指令优化的关键在于深入理解目标硬件架构特性，并将其映射到算法实现中，例如重新排列数据以适应 SIMD 架构、调整算法以利用特定硬件加速器。此外还包括编译器级别优化：自动向量化、指令调度、寄存器分配等。在本系列 05OtherOpt.md 中，还详细介绍了张量化（Tensorization）——如mma.sync.aligned.m8n8k32.row.col.load.128b这类 Tensor Core 指令可在一条指令内完成 $C[8\times8] += A[8\times32] \times B[32\times8]$ 的矩阵乘累加，以及通过拆分算子将部分计算映射为硬件微内核（如 TVM 的 tensorize）的思路。

存储优化：延迟隐藏与内存层次利用

存储优化通过精细调整数据访问模式和内存层次结构提升整体性能：

• 内存延迟隐藏：通过预取（Prefetching）数据到缓存、调整数据访问模式以提高缓存命中率，或并行执行其他计算任务，让处理器在等待数据加载时保持忙碌，提高资源利用率。
• 双缓冲（Double Buffering）：使用两个缓冲区平滑数据流、隐藏延迟——一个缓冲区的数据正在处理时，另一个缓冲区可加载新数据，特别适用于图形渲染、视频处理等连续数据流场景。
• 资源管理：合理分配内存资源、优化数据结构以减少内存占用、使用内存池减少内存分配/释放开销。

这些策略共同作用可显著提高内存访问效率，减少因内存瓶颈导致的性能损失。

DSL 开发算子：更高生产效率的折中

手写算子的开发往往要求深入底层硬件细节，涉及数据布局、指令选择、索引计算等精细调整，显著提高 Kernel 编写难度。为简化这一过程，业界发展出多种领域特定语言（DSL），其中TVM与Triton是杰出代表。

这些 DSL 通过高级抽象封装常用优化技术，并由编译器的优化阶段自动识别与应用。开发者只需专注高层次计算逻辑，利用 DSL 的 API 即可实现高性能算子。相比手写代码和极限优化，这种方法虽然可能无法达到极致性能，但通常能实现超过 90% 的性能水平，同时开发效率提升数十倍——这种权衡在多数场景下是合理的。

TVM：计算与调度分离的端到端 AI 编译器

TVM 是 OctoML 研发的可扩展、开放的端到端 AI 编译器，用于神经网络模型的优化和部署，目标是减少企业为特定硬件开发和深度学习软件部署所花费的成本和时间。

TVM 极大地发扬了 Halide 的计算与调度分离思想，将可实现的优化都以调度 API 的形式呈现（继承自 Halide 的调度原语包括 Reorder、Split、Fuse、Tile、Vectorize、Unroll、Parallel 等，详见 02OPSCompute.md）。相比 Triton，TVM 不局限于 CUDA，支持更多后端，也易于扩展新后端。TVM 上一个典型的调度如下（来自 VTA 示例中的 packed dense 调度）：

def schedule_dense_packed(cfg, outs): """Packed dense schedule.""" assert len(outs) == 1 output = outs[0] const_ops = [] ewise_inputs = [] ewise_ops = [] dense_res = [] assert "int" in output.op.input_tensors[0].dtype def _traverse(op): if topi.tag.is_broadcast(op.tag): if not op.same_as(output.op): if not op.axis: const_ops.append(op) else: ewise_ops.append(op) for tensor in op.input_tensors: if isinstance(tensor.op, tvm.te.PlaceholderOp): ewise_inputs.append((op, tensor)) else: _traverse(tensor.op) else: assert op.tag == "dense_pack" dense_res.append(op) _traverse(output.op) assert len(dense_res) == 1 dense_stage = dense_res[0].output(0) s = te.create_schedule(output.op) ##### space definition begin ##### b, c_o, _, _ = s[dense_stage].op.axis c_i, _ = s[dense_stage].op.reduce_axis cfg.define_split("tile_b", b, num_outputs=2) cfg.define_split("tile_ci", c_i, num_outputs=2) cfg.define_split("tile_co", c_o, num_outputs=2) cfg.define_knob("oc_nthread", [1, 2]) ###### space definition end ###### data, weight = dense_stage.op.input_tensors env = get_env() cdata = s.cache_read(data, env.inp_scope, [dense_stage]) cweight = s.cache_read(weight, env.wgt_scope, [dense_stage]) s[dense_stage].set_scope(env.acc_scope) # cache read input cache_read_ewise = [] for consumer, tensor in ewise_inputs: cache_read_ewise.append(s.cache_read(tensor, env.acc_scope, [consumer])) # set ewise scope for op in ewise_ops: s[op].set_scope(env.acc_scope) s[op].pragma(s[op].op.axis[0], env.alu) for op in const_ops: s[op].compute_inline() # apply tiling for SRAM reuse x_b, x_c, _, _ = s[output].op.axis x_bo, x_bi = cfg["tile_b"].apply(s, output, x_b) x_co, x_ci = cfg["tile_co"].apply(s, output, x_c) s[output].reorder(x_bo, x_co, x_bi, x_ci) store_pt = x_co # set all compute scopes s[dense_stage].compute_at(s[output], store_pt) for op in ewise_ops: s[op].compute_at(s[output], store_pt) for tensor in cache_read_ewise: s[tensor].compute_at(s[output], store_pt) s[tensor].pragma(s[tensor].op.axis[0], env.dma_copy) # virtual threading along output channel axes if cfg["oc_nthread"].val > 1: _, v_t = s[output].split(x_co, factor=cfg["oc_nthread"].val) s[output].reorder(v_t, x_bo) s[output].bind(v_t, te.thread_axis("cthread")) x_bo, x_co, x_bi, _ = s[dense_stage].op.axis k_o, _ = s[dense_stage].op.reduce_axis s[dense_stage].reorder(x_bo, k_o, x_co) k_o, _ = cfg["tile_ci"].apply(s, dense_stage, k_o) s[cdata].compute_at(s[dense_stage], k_o) s[cweight].compute_at(s[dense_stage], k_o) # Use VTA instructions s[cdata].pragma(s[cdata].op.axis[0], env.dma_copy) s[cweight].pragma(s[cweight].op.axis[0], env.dma_copy) s[dense_stage].tensorize(x_bi, env.gemm) s[output].pragma(x_ci, env.dma_copy) return s

这段调度代码展示了 TVM 将多种优化手段组合为调度空间的典型范式，可拆解为几个关键层次：

• 调度空间定义：通过cfg.define_split对 batch（tile_b）、reduction 轴（tile_ci）和输出通道（tile_co）定义可搜索的拆分选项，通过cfg.define_knob("oc_nthread", [1, 2])定义虚拟线程数，为后续自动调优器提供搜索维度。
• 存储/作用域设置：cache_read将输入缓存到 VTA 的输入（inp_scope）和权重（wgt_scope）片上存储，set_scope将 dense 累加器与逐元素算子放到加速器累加作用域，常量算子compute_inline内联。
• 循环变换：对输出轴做分块（tiling）后reorder重排，以提升 SRAM 复用；compute_at将 dense 阶段、逐元素算子和缓存读取的计算挂载到输出循环的存储点（store_pt）下，形成合理的计算顺序与作用域嵌套。
• 指令映射与张量化：通过pragma(…, env.dma_copy)将数据搬运映射为 DMA 指令，通过tensorize(x_bi, env.gemm)将最内层循环替换为 VTA 的 GEMM 硬件指令，pragma(x_ci, env.dma_copy)将输出写回映射为 DMA 拷贝——这就是将"循环优化 + 存储优化 + 指令优化"三者统一编码进调度树的实例。

使用 TVM 提供的调度 API，可以轻松实现循环、指令、存储等层面的优化，还可以将这些优化参数化，利用自动调优器搜索更高效的实现（自动调优原理可参考本系列 06AutoTuning.md）。

TVM 当前被多个加速器厂商采用，进行了适应自家硬件的定制开发：

• 希姆计算：基于 TVM 的 AI 编译器端到端支持希姆一代、二代芯片，实现了自定义算子方案，其模型性能接近手写极致。
• 华为 TBE 张量加速引擎：TBE（Tensor Boost Engine）负责执行昇腾 AI 处理器中运行在 AI Core 上的算子，TBE 提供了基于 TVM 框架的自定义算子开发能力，通过 TBE 提供的 API 可以完成相应神经网络算子的开发。

Triton：基于分块编程范式的 GPU 编程语言

Triton 是 OpenAI 研发的专为深度学习和高性能计算任务设计的编程语言和编译器，旨在简化并优化在 GPU 上执行的复杂操作的开发。其目标是提供开源环境，以比 CUDA 更高的生产力编写快速代码。

Triton 的核心理念是：基于分块的编程范式可以有效促进神经网络高性能计算核心的构建。CUDA 的编程模型是传统的 SIMT（Single Instruction Multiple Thread）GPU 执行模型，在线程的细粒度上进行编程；Triton 则是在分块（Block）的粒度上进行编程。例如在矩阵乘法场景下，CUDA 与 Triton 存在如下差异：CUDA 需要程序员在 Thread 级别手工管理每个线程负责的元素与访存模式，而 Triton 在循环中是逐块进行计算的。

这种方法的一个关键优势是：块结构的迭代空间相比现有 DSL，为程序员实现稀疏操作提供了更多灵活性，同时允许编译器为数据局部性和并行性进行积极优化。下面是一个使用 Triton 实现矩阵乘法的完整示例：

@triton.jit def matmul_kernel( a_ptr, b_ptr, c_ptr, stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn, M: tl.constexpr, N: tl.constexpr, K: tl.constexpr, # M=N=K=1024 BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr, #BLOCK_M=BLOCK_N=BLOCK_K=32 ): offs_m = tl.arange(0, BLOCK_M) offs_n = tl.arange(0, BLOCK_N) offs_k = tl.arange(0, BLOCK_K) a_ptrs = a_ptr + offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak b_ptrs = b_ptr + offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn accumulator = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32) for k in range(0, K, BLOCK_K): a = tl.load(a_ptrs) b = tl.load(b_ptrs) accumulator += tl.dot(a, b) a_ptrs += BLOCK_K * stride_ak b_ptrs += BLOCK_K * stride_bk c_ptrs = c_ptr + offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn tl.store(c_ptrs, accumulator)

这段代码中，开发者只声明了块内索引（offs_m/offs_n/offs_k）与分块大小（BLOCK_M/BLOCK_N/BLOCK_K，示例中取 32），通过tl.arange构造块内坐标、tl.load/tl.store完成块级访存、tl.dot完成块级矩阵乘，K 维度上的累加由外层for k in range(0, K, BLOCK_K)驱动——线程划分、共享内存管理、数据搬运等底层细节全部交由 Triton 编译器处理。

在应用场景上，Triton 已被集成进多个著名项目中：

• jax-triton（JAX）：JAX 是用于加速数值计算的 Python 库，使用 Triton 编写可嵌入 JAX 程序的自定义 GPU 内核，在 JAX 中可通过triton_call方便地调用 Triton kernel。
• PyTorch/Inductor：Inductor 在 Triton 集成方面做得更加全面且务实，共包含三种使用方式：针对非计算密集算子，基于 Inductor IR 实现相对通用的 Codegen 支持；针对 GEMM，基于 Jinja2 通过模式匹配实现半定制化的 codegen；针对 Conv，调用预置（pre-baked）的 Triton kernel，不提供定制能力。

Triton 实现原理：基于 MLIR 的编译器栈

手写 CUDA Kernel 的三大挑战

在没有 Triton 之前，算子工程师开发算子时需要同时处理 DRAM、SRAM、计算单元，面临诸多挑战：

• 内存管理：合理利用内存体系，将频繁访问的数据块缓存到较快的存储区域；对齐和合并访存请求，避免带宽浪费。
• 线程管理：最大化利用硬件计算资源，规划并行线程数量和线程束（warp）大小。
• 指令使用：使用 CUDA 实现同一功能有多种指令可选，不同指令具有不同延迟和吞吐量。

Triton 提高了算子开发效率，使开发者不再囿于硬件细节。CUDA 直接面向 Thread 编程，而 Triton 面向Thread Block编程，开发者只需关注三点：1）Kernel launch 的参数；2）每个数据分块的大小；3）数据分块之间的交互。这之下的细节由 Triton 实现。

Triton 整体架构

Triton 是基于 MLIR 实现的，其架构分为前端（Frontend）与优化器（Optimizer）两大层次：

Frontend负责将开发者用 Python 编写的 kernel 转换为对应的 Triton IR（Triton Dialect）。使用@triton.jit标注 kernel 后，Triton 解析 Python AST，将用户定义的计算过程带入 MLIR 体系，之后继续做后续优化。

从 TritonIR 到 TritonGPU IR 的优化管线

Optimizer 的大致工作流如下：

主要分为三个阶段：1）TritonIR 的优化；2）TritonIR 到 TritonGPU IR 的转换；3）TritonGPU IR 的优化。贯穿中间的数据结构是TritonGPU IR。

TritonGPU Dialect相比 Triton Dialect，主要增加了 GPU 硬件相关的 Op 和 Type，关键 Op 为数据布局转换。当前有以下几种数据布局：

• Blocked Layout：表示 thread 间平均分配 workload 的情况，每个线程处理一块内存上连续的数据。
• Shared Layout：表示数据在 shared memory 中的一些特性。
• MMA Layout：表示 Tensor Core 中 MMA 指令结果的 data layout。
• DotOperand Layout：表示 Triton 的 DotOp 输入的 layout。

这些 layout 信息会作为 attribute 附着在 Tensor 对象上，描述该 Tensor 作为输入或输出时所需满足的映射关系；若出现输入与输出映射关系不兼容的情况，编译器会插入中转 layout 完成兼容性适配（代价是可能引入额外转换开销）。典型的 layout 转换及其特点：

• #shared -> #blocked：数据从 shared memory 加载到 register file，需要考虑 swizzle（避免 bank conflict）。
• #blocked -> #shared：数据从 register file 存储到 shared memory，需要与上一步相同的 swizzle 方式。
• #mma -> #blocked：DotOp 的输出转换为更简单的 layout 以进一步计算，由于涉及跨 thread 数据传递，一般借由 shared memory 中转一次。
• #blocked -> #dot_operand：转换为 DotOp 的输入，这一步可能也需要 shared memory 中转。

TritonIR 上的优化主要是与硬件无关的计算优化，包含如下 Pass：1）Inliner Pass，将 Kernel Call 的子函数内联展开；2）Combine Pass，执行一些特定 Pattern 的 rewrite；3）Canonicalizer Pass，执行化简类的 Pattern rewrite；4）CSE Pass，消除公共子表达式；5）LICM Pass，将循环无关变量移到 for 循环外。

TritonGPU IR 优化在计算优化之外，新增了 GPU 硬件相关优化，具体 Pass 如下：1）ConvertTritonToTritonGPU Pass，将 Triton IR 转换为 TritonGPU IR，主要是增加 TritonGPU 特有的 layout；2）Coalesce Pass，重排 order，使最大连续性（contiguity）的维度排在最前面，辅助向量化访存；3）Pipeline Pass，MMA 指令对应的 global memory 到 shared memory 的 N-Buffer 优化，缓解计算与访存差异；4）Prefetch Pass，MMA 指令对应的 shared memory 到 register file 的 N-Buffer 优化。

综上，用户在开发 Kernel 时主要关注业务逻辑，底层硬件优化细节由 Triton 编译器实现；但对于一些十分精细的优化，使用 Triton 可能就无法实现，仍需回到手写 Kernel 路线。

小结与思考

本节围绕"算子手工优化"形成完整的方法论闭环：

• 计算分析：通过定义计算量、访存量、计算强度和理论性能等指标，利用 Roofline 模型分析程序瓶颈（Compute-Bound / Memory-Bound），指导优化策略的制定。典型案例是 MobileNet 在 1080Ti 上受带宽限制、VGG16 可完全利用算力。
• 优化策略：针对算子的计算特性和硬件资源特点，采用循环优化（分块、展开、重排、融合、拆分）、指令优化（向量化、张量化）、存储优化（延迟隐藏、双缓冲、预取）等技术提升性能；仓库 04Backend/code 提供了循环优化四类变换的可运行伪代码佐证。
• DSL 开发算子：使用 TVM（调度 API + 自动调优）和 Triton（分块编程 + MLIR 优化管线）等高级语言加速算子开发，通过封装优化技术和自动调优，以 90% 以上的性能换取数十倍的开发效率提升。

延伸阅读

本系列更多相关内容，可继续阅读同一目录下的姊妹章节：02OPSCompute.md（计算与调度、调度树与调度变换）、04LoopOpt.md（循环优化细节与 Cache 局部性分析）、05OtherOpt.md（指令与存储优化）、06AutoTuning.md（自动调优原理）；配套课件与视频资源见 03Optimization.pdf 与 03Optimization.pptx。Triton MLIR 架构与优化流程部分的整理，参考了《OpenAI/Triton MLIR 迁移工作简介》公开技术博客。

• 文档
• 教程
• 人工智能