AI 编译器与 Kernel 编程（Triton/Pallas/cutlass）

• 理解 StableHLO 的定位与来源
• 掌握统一 IR 对跨框架的价值
• 理解 StableHLO 到 MLIR/LLVM 的 lowering

StableHLO 是 Google 从 XLA 的高层 HLO 拆分出的稳定、可版本化的中间表示（IR），作为 ML 框架（PyTorch、JAX、TensorFlow）与硬件后端之间的统一层。各框架把计算图降为 StableHLO，StableHLO 再通过 MLIR 逐步 lowering 到 GPU/LLVM 代码。工程价值：统一 IR 让不同框架共享同一套编译后端，避免为每个框架单独实现后端；StableHLO 的稳定 ABI 与版本化保证框架间、前端与后端解耦，便于扩展新硬件；MLIR 作为工具链提供多级 lowering 与优化（算子融合、布局、代码生成），降低新框架/新硬件接入成本。一句话，StableHLO 提供"一次编译、多框架复用、多后端适配"的中间层，是跨框架编译的基石。

StableHLO 的价值在于"统一"与"稳定"：统一 IR 让前端（框架）与后端（硬件）解耦，稳定版本化让生态可维护。理解其作为中间层的角色，就能理解跨框架编译器为何都围绕它构建。

• 理解 IREE 的定位
• 掌握多目标编译能力
• 理解从 StableHLO 到各后端的流程

IREE 是一个基于 MLIR 的编译器与运行时，以 StableHLO（及其他前端 IR）为输入，编译到多种后端目标：Vulkan（跨 GPU 厂商的图形计算 API）、CUDA（NVIDIA GPU）、CPU（LLVM 生成）、以及 WebGPU 等。工程价值：高度可移植——同一份 StableHLO 计算图可编译并部署到 Vulkan/CUDA/CPU/Web 等平台，覆盖从数据中心到移动端/浏览器的场景；IREE 提供模块化、可扩展的编译与运行时，支持异步执行、内存管理与流式调度，便于在异构设备上部署推理。多目标价值的核心是"一次编译、多平台部署"，降低在多种硬件上维护推理栈的成本。

IREE 的价值是"多目标编译 + 可移植运行时"。理解它从 StableHLO 统一 IR 出发、面向 Vulkan/CUDA/CPU/Web 多后端，就能把握跨平台推理部署的思路。

• 理解 XLA 的编译流程
• 掌握 fusion、tile、bufferize 的作用
• 理解对 TF/JAX 的性能价值

XLA 是 TF/JAX 的编译加速器，核心流程包括：high-level op fusion（高层算子融合），把多个算子合并为单个（如融合点乘、加、激活、归一化），减少 kernel launch 与中间张量访存；tile（分块），把计算划分为适合硬件缓存与并行执行的分块，提升数据复用与并行度；bufferize（缓冲化），把无副作用的张量计算转换为显式内存分配与复用，决定中间数据的落点与生命周期，减少内存压力。工程价值：XLA 把 TF/JAX 的优化 / eager 计算编译为高效 GPU/TPU 代码，使研究者无需手写 kernel 即可获得接近手工优化的性能，并统一了 JIT 编译路径。fusion+tile+bufferize 的组合是 XLA 提速的核心，也是其能在 TPU/GPU 上高效运行的保证。

XLA 的编译价值在"高层融合 + 分块 + 内存缓冲化"三件套。fusion 减访存、tile 提并行、bufferize 管内存，理解这三步就理解了 XLA 的编译主线。

• 理解 cuBLAS 与 Triton/cutlass 的层次
• 掌握 Triton 与 cutlass 的差异
• 理解替换的适用边界

cuBLAS 提供预编译、高度优化的常用矩阵运算，性能好但不可定制或替换。Triton 与 cutlass 是替代/定制方案：Triton 用 Python 块级 DSL 编写 kernel，编译器自动生成 CUDA 代码，开发效率高、易调整算法，但抽象层失去底层控制（寄存器分配、精确调度），性能可能不如手工优化；cutlass 是 C++ 模板库，提供精细控制（tile 大小、warp 布局、TMA、WGMMA 等），能用底层指令实现接近理论上限的性能，但开发复杂、需要深度硬件知识。替换边界：当 cuBLAS 无法满足需求（自定义算子、融合、特殊稀疏/形状、或需要结合特定硬件特性）时，用 Triton 快速迭代或 cutlass 深度优化。cuBLAS 仍是默认选择，Triton/cutlass 用于 cuBLAS 覆盖不佳或追求极致性能的场景。

cuBLAS 是"拿来即用"，Triton 是"易写易改"，cutlass 是"极致可控"。替换边界取决于性能需求与定制程度：常规用 cuBLAS，需定制或极致性能时选 Triton/cutlass。

• 理解 Triton 的块级抽象
• 掌握 Python DSL 的工程价值
• 理解 matmul/attention 的编写

Triton 用 Python DSL 编写块级（block-level）GPU kernel，程序员以块/tile 为单位描述运算，无需手动管理线程块、共享内存与寄存器分配。工程价值：开发效率高——几十行 Python 即可写出高性能 matmul/attention kernel，编译器自动处理 tiling、数据布局与内存流水；可读性好——块级表达贴近数学公式，便于理解和维护；可移植性强——同一 Triton 代码可编译到不同 GPU（CUDA/ROCm）。在 matmul/attention 上，Triton 通过 @triton.jit 装饰 kernel，用 tl.load/store 加载块、用块级张量运算实现分块矩阵乘与注意力，配合 torch.compile 自动生成 Triton kernel。工程价值在于让"写 kernel"从底层手工优化变成"高层描述 + 编译器优化"，大幅降低高性能算子开发门槛。

Triton 的工程价值是"块级抽象 + 编译器自动化"。它把 tiling、内存、调度交给编译器，程序员专注算法，从而在易用性与性能间取得平衡。是 PyTorch 2.x 编译栈的核心。

• 理解 Inductor 的角色
• 掌握 Triton 代码生成
• 理解对 torch.compile 的价值

TorchInductor 是 PyTorch 2.x 中 torch.compile 的默认后端，负责把捕获的计算图编译为高效代码。在 GPU 上，Inductor 通过 codegen 生成 Triton kernel：把图优化（算子融合、layout 变换）后的算子映射为 Triton 块级代码，交给 Triton 编译为 CUDA。工程价值：Inductor 把"图优化"与"kernel 生成"结合，让 torch.compile 能自动把 eager 计算转换为高性能 Triton kernel，无需手写；Inductor 处理循环、融合、内存访问与 autotune，使同一模型在 GPU 上获得接近手工优化的性能。它把 PyTorch 的编译能力与 Triton 的 kernel 生成打通，是 torch.compile 性能提升的关键组件，也是扩展新后端（如 ROCm）的入口。

Inductor 是 torch.compile 的"代码生成器"，在 GPU 上用 Triton 生成 kernel。理解 Inductor 的 codegen 价值，就能理解 torch.compile 为何能自动提升性能并支持多后端。

• 理解 Relax 的定位
• 掌握对 dynamic shape 的支持
• 理解 control flow 的处理

TVM Relax 是 Relay 的继任者，用于处理动态形状（dynamic shape）与控制流（control flow）的图 IR。Relay 在处理动态 shape 与复杂控制流（如 if/while、可变循环）时受限，Relax 通过更灵活的 IR 设计支持运行时形状未知、带有分支与循环的图，并把这些高层结构降到可执行的代码。工程价值：Relax 让编译器能处理生产场景中常见的动态输入（如变长序列、动态 batch）与复杂控制流（如条件解码、循环），而非只支持静态 shape 的简单图；同时 Relax 与 TVM 的 TensorIR/调度结合，支持在这些动态结构下做优化与代码生成。它拓宽了 TVM 的适用范围，使其能编译更真实、更动态的模型。

Relax 的价值在于扩展 IR 的表达能力（动态 shape、控制流）。理解它比 Relay 更灵活，就能理解 TVM 如何应对真实生产中的动态模型。

• 理解 tile-aware autotune 机制
• 掌握 persistent kernel 概念
• 理解两者对性能的价值

Triton 的 tile-aware autotune 通过 @triton.autotune 为 kernel 提供多个候选配置（tile 大小、num_warps、num_stages），在运行时基准测试选择最优并缓存，从而针对不同硬件与输入自动选择最佳 tile 参数，提升性能与可移植性。persistent kernel（持久化 kernel）指一个 kernel 启动后长期驻留，通过循环处理大量数据块，避免反复启动 kernel 的开销；Triton 支持在 kernel 内用循环处理多次 tile，结合 persistent 模式减少启动开销、提高流水线效率。工程价值：tile-aware autotune 让同一 kernel 在不同硬件上自动适配最优 tile，persistent kernel 减少启动开销、提升吞吐，两者结合让 Triton 写出高性能且可移植的 kernel。理解它们能指导在 Triton 中做性能调优。

autotune 解决"选对 tile"，persistent kernel 解决"减少启动"。两者都是 Triton 获得高性能的关键手段，理解了它们就理解了 Triton 性能调优的两个方向。

• 理解 StableHLO 与 MHLO 的关系
• 掌握兼容性设计
• 理解迁移价值

StableHLO 与 MHLO 都是 XLA 高层 HLO 的 MLIR 方言。MHLO 是较早的 XLA HLO 的 MLIR 表示，StableHLO 是后来从 MHLO 拆分、专门为稳定与可版本化设计的方言。工程价值：StableHLO 提供稳定的 ABI 与标准化方言，保证跨版本、跨框架的兼容性，避免 MHLO 随 XLA 内部演进而发生变动；StableHLO 与 MHLO 以 module 为单位，通过兼容性转换（如 StableHLO→MHLO 的 lowering 或兼容检查）协同，使基于 MHLO 的工具链能渐进迁移到 StableHLO。这种兼容性设计让上游框架（JAX/TF/PyTorch）与下游编译器解耦，生态更稳定、可维护。StableHLO 作为"稳定接口"的价值在于减少 ML 生态的版本破碎风险。

StableHLO 与 MHLO 的兼容性价值是"稳定接口 + 平滑迁移"。StableHLO 提供稳定 ABI，MHLO 作为早期实现，二者通过 module 级兼容转换协同，降低生态破碎风险。

• 理解 IREE 的多平台部署能力
• 掌握各平台的编译后端
• 理解适用边界

IREE 支持多种平台部署：Android 通过 Vulkan/CPU 后端运行模型，Web 通过 WebGPU/WebAssembly 让模型在浏览器运行，GPU 通过 Vulkan/CUDA 部署。工程边界：IREE 适合需要跨平台、可移植部署的推理场景（边缘设备、移动端、浏览器），其模块化运行时与多后端支持让同一模型编译到多平台；但 IREE 的边界在于——对特定平台/硬件的极致性能优化能力不如厂商原生方案，且新算子/新硬件支持依赖其生态与 lowering 覆盖。IREE 的价值在于"一次编译、多平台部署"，适用于跨 Android/Web/GPU 的统一部署；若追求单一平台的极致性能，可能需要结合原生后端。

IREE 的边界是"可移植性 vs 极致性能"。它擅长跨平台统一部署，但对特定平台的深度优化能力有限。理解边界才能选对部署方案。

• 理解 meta schedule 的概念
• 掌握 auto-tuning 的机制
• 理解工程价值

TVM 的 meta schedule（Meta 调度）是自动搜索调度策略的框架，通过搜索空间设计、成本模型与进化/采样器，自动为算子找到最优调度（tile 大小、循环顺序、向量化等）。auto-tuning（自动调优）在真实硬件上运行候选调度并测量性能，用成本模型预测 + 实测反馈迭代优化，找到最优实现。工程价值：meta schedule + auto-tuning 让编译器无需人工手工调优，便能针对不同硬件（GPU/CPU/TPU）自动生成高性能 kernel，显著降低性能迁移成本；同时可持续改进（自适应搜索），把工程经验沉淀为搜索策略。它把"性能工程师的手工调优"自动化，是 TVM 相比手写 kernel 的关键优势。

meta schedule 与 auto-tuning 的价值是"把性能调优自动化"。理解搜索空间、成本模型与实测反馈，就能理解 TVM 如何自动生成高性能 kernel。

• 理解 TensorIR 的定位
• 掌握 tensor-level 代码生成
• 理解调度与生成的结合

TVM TensorIR 是用于张量级中间表示与调度的 IR，把张量算子表示为可调度的循环结构，支持 tiling、向量化、并行化、拆分/合并、数据布局等调度变换，然后把调度后的循环生成目标代码（CUDA/ROCm/LLVM）。工程价值：TensorIR 让 kernel 开发以"张量级 + 调度"进行，而非手写底层汇编；程序员描述算子与调度意图，编译器生成代码，调度可复用、可搜索、可自动优化。作为 tensor-level primitive 的 codegen，TensorIR 与 meta schedule 结合，能自动搜索最优调度并生成高性能 kernel，是 TVM 生成高性能算子的核心。它的价值在于把"算子实现"与"调度优化"分离，兼顾表达力与自动化。

TensorIR 的价值是"张量级调度 + 代码生成"。它分离了算子语义与调度实现，让调度可自动搜索，是 TVM 高性能 codegen 的基础。

• 理解 TVM Unity 的定位
• 掌握统一编译流水线
• 理解与其他组件的协同

TVM Unity 是 TVM 新架构，把编译流水线统一化：通过模块化、可组合的 IR（Relax、TensorIR、TIR）与统一编译流程，让前端（图级 IR）、调度（TensorIR）与代码生成（后端）在同一框架内协同，使编译过程可插拔、可扩展。工程边界：TVM Unity 主要的价值是统一并简化了 TVM 的编译流水线，让不同前端（Relax/Relay）与后端（CUDA/ROCm/CPU）共享一致的编译与优化流程，减少重复；但它的边界在于——统一流水线覆盖的算子/后端有一定范围，前沿硬件或新算子可能需要扩展 lowering 与成本模型。统一协同的价值是降低维护成本、让优化可复用，边界是生态覆盖与深度优化范围。

TVM Unity 的价值是"统一 + 可组合的编译流水线"。理解它让图级/张量级/后端 IR 协同，能把握 TVM 新架构的演进方向与适用边界。

• 理解 torch.compile 的编译流程
• 掌握 eager 到 inductor 的转换
• 理解工程价值

torch.compile 把 PyTorch 2.x 的 eager mode（逐算子执行）计算捕获为图表，经 aot_autograd 生成反向图，再由 Inductor backend 做 codegen 生成 Triton/C++ 代码。工程价值：eager mode 灵活但逐算子调度（Python 解释、kernel launch）开销大；torch.compile 通过图捕获 + 编译，消除逐算子调度开销，利用算子融合、代码生成与 autotune 提升性能，同时保持 eager 的便利（通过装饰器或 torch.compile API 一键启用）。它为研究者提供"无需改代码即可加速"的路径，让 eager 的灵活性与编译的高性能结合，是 PyTorch 性能提升的关键。工程价值在于把"手写 kernel 优化"变成"自动编译优化"。

torch.compile 的价值是"eager 的便利 + 编译的性能"。理解 eager→图捕获→Inductor codegen 的流程，就理解它如何自动加速 PyTorch 模型。

• 理解各编译组件的角色
• 掌握 graph capture 流程
• 理解分工价值

torch.compile 的编译流程含多个组件：aot_eager 是把图用 eager 语义先执行/捕获，用于验证与调试；aot_autograd 通过 AOT 高阶自动微分，同时产生前向与反向图并用编译优化；inductor 作为后端把图编译为 Triton/C++ 代码。在 graph capture 模式下，optimized eager 先捕获计算图，aot_autograd 生成可编译的反向图，inductor 做融合与 codegen。工程价值：分层让各组件职责清晰——aot_eager 保证正确性、aot_autograd 处理反向与图优化、inductor 负责代码生成，从而支持渐进式编译（从 eager 到 fully compiled）、可调试与可扩展。理解这些组件能把握 torch.compile 的编译流水线与各阶段的优化点。

aot_eager/aot_autograd/inductor 是 torch.compile 流水线的三个环节：捕获、反向+图优化、代码生成。理解分工才能调试与优化编译流程。

• 理解 WGMMA 指令
• 掌握 4-warp group MMA 机制
• 理解工程价值

WGMMA（Warpgroup Matrix-Multiply-Accumulate）是 Hopper（SM90）引入的指令，让一个 warpgroup（由 4 个 warp 组成，128 线程）协同执行大块矩阵乘加，比之前按 warp 的 MMA 指令更高效。工程价值：WGMMA 以 warpgroup 为单位做大块 MMA，减少指令发射与调度开销、提升 Tensor Core 利用率；配合 TMA（异步批量加载）与 shared memory 流水线，可让数据加载与 MMA 计算重叠，把矩阵乘推到带宽/算力上限。Blackwell 延续并强化这类指令。WGMMA 的工程价值在于提供"大块、协同、硬件原生"的矩阵运算路径，让 GEMM 类算子（如 cutlass 实现）能更高效地利用 Tensor Core。

WGMMA 的价值是"更大粒度、跨 warp 协同的 MMA"。它配合 TMA 做流水线，是 Hopper/Blackwell 上高性能 GEMM 的关键指令。理解 warpgroup 粒度就能理解其性能优势。

• 理解依赖分析
• 掌握 tiling 的合法性条件
• 理解局部性提升

loop tiling 把循环迭代域划分为小块（tile），先执行完一个 tile 再执行下一个，从而提升数据在缓存中的复用（局部性）。合法性由依赖分析保证：tiling 必须不破坏数据依赖。若循环中某次迭代读写的数据被另一迭代使用，tiling 改变迭代执行顺序，需确保依赖关系（距离向量）不被破坏。具体地，若依赖距离向量在 tiled 维度上方向一致且 tile 边界不切断依赖（即依赖的迭代落在同一或相邻 tile 内且顺序保持），则 tiling 合法。对可并行循环（无跨迭代依赖或依赖粒度为 tile 内），tiling 合法。局部性提升源于：tile 内的数据在一次访存中被多次使用（时间局部性），且同一 tile 的数据在空间上相邻（空间局部性），从而减少缓存 miss 与全局内存访问。tiling 是编译器优化中兼顾缓存与并行的重要手段。

tiling 的合法性依赖分析：依赖不被 tile 边界切断。局部性提升来自 tile 内数据复用。理解依赖分析是判断 tiling 是否合法、何时可自动化的前提。

• 理解多面体模型
• 掌握迭代域与访问的线性表示
• 理解统一调度变换

多面体模型（polyhedral model）用整数线性约束（仿射不等式）把循环 nest 的迭代域表示为多面体（polyhedron）：迭代向量（i,j,k）满足一系列仿射不等式即构成迭代域。数组访问（如 A[i][j+k]）也用仿射函数表示，从而可以精确刻画访问关系与数据依赖。调度（schedule）是迭代向量到执行顺序的仿射映射，通过改变调度（如 tiling、interchange、skewing）可统一地变换循环结构，而依赖分析在多面体上通过比较两个访问的仿射表达式判断依赖。统一调度变换的价值：在统一的数学框架（多面体 + 仿射调度）下，所有循环变换（分块、交换、分片、并行化）都是对调度矩阵的修改，自动且可验证地保持合法性，并同时优化局部性与并行性。这是多面体编译（如 Pluto、TensorIR 的依赖）与并行循环优化的理论基础。

多面体模型把循环问题转化为"多面体 + 仿射调度"的数学问题。迭代域用仿射约束表示、访问用仿射函数表示、变换用调度映射表示，从而统一且自动化地做优化。

• 理解 tile 与缓存容量的关系
• 掌握 tile 过大/过小的损失
• 理解容量匹配原则

tile size 的选择要与目标缓存层次（如 L1/L2 或共享内存）容量匹配，因为 tile 内数据需要在缓存中复用，若 tile 过大导致数据超过缓存容量，会触发缓存溢出（capacity miss），数据被反复换出/换入，局部性收益丧失，甚至退化为类似无 tiling 的访存；若 tile 过小，则 tile 内数据复用率低，每次 tile 都要重新加载数据，缓存命中受益有限，且 tile 切换/循环开销增加。因此最优 tile size 是"缓存容量内尽可能大"：既充分利用缓存复用数据，又不超出容量导致 miss。在 GPU 上，tile 大小还受共享内存容量与寄存器数约束，过大则放不下，过小则复用不足。匹配容量的原则是让 tile 的数据集在缓存/SRAM 中驻留，最大化复用同时避免溢出。

tile 大小是缓存容量与复用率的权衡。过大→capacity miss，过小→复用不足。理解"缓存容量内最大"即可把握 tile 大小的选择原则。

• 理解仿射要求的原因
• 掌握非仿射访问的处理
• 理解抽象解释

多面体编译要求循环边界与数组下标是仿射（线性）函数，因为多面体模型的数学基础（迭代域表示为仿射不等式组、访问关系表示为仿射函数、依赖通过仿射比较判定）依赖线性结构。只有仿射边界与下标，迭代域才是多面体、访问关系才能精确分析、依赖与调度变换才能合法地自动推导。若边界或下标非仿射（如循环边界依赖运行时数据、下标含 x[i*idx[i]] 这类间接寻址），多面体模型无法精确表示，依赖分析会丢失或出错。非仿射访问的处理：保守分析（把非仿射访问视为可能访问整个数组，取保守依赖，虽可能过度保守但保证正确）、或把部分非仿射信息剥离/抽象化、或放弃多面体优化而用手工/通用优化。总之，非仿射访问使多面体方法的精确性下降，需保守处理或降级为其他优化。

仿射要求是多面体模型数学自洽的前提。非仿射访问破坏了精确性，需保守抽象或降级。理解这一点就理解了多面体编译的适用边界。

• 理解 cutlass 3.x 与 cuBLASLt 的定位
• 掌握性能协同点
• 理解工程边界

cuBLASLt 是 NVIDIA 的轻量级矩阵运算库，提供 GEMM 的启发式/自动调优与多种布局，开箱即用、性能好；cutlass 3.x 是 C++ 模板库，提供底层控制（tile、warp 布局、TMA、WGMMA）与自定义 kernel 开发。性能协同：cutlass 实现的高性能 GEMM 常为 cuBLASLt 提供优化的参考实现，二者共享底层 Tensor Core 优化思路；开发者可用 cuBLASLt 快速获得良好性能，再用 cutlass 深度定制/替换瓶颈。工程边界：cuBLASLt 是黑盒、不可定制但省心；cutlass 是白盒、可深度优化但需专业知识。协同的价值是"先用 cuBLASLt 保底，再用 cutlass 突破"，边界是性能需求与定制深度——常规用 cuBLASLt，追求极致/定制用 cutlass 并可与 cuBLASLt 协同。

cuBLASLt 与 cutlass 的协同是"黑盒保底 + 白盒突破"。cuBLASLt 省心、cutlass 深度可控，边界取决于性能需求与定制深度。

• 理解 Pallas 与 pallas_call
• 掌握 TPU low-level 编写
• 理解工程价值

Pallas 是 JAX 的低层 kernel 编程接口，通过 pallas_call 让用户在 TPU 上直接编写 low-level kernel，提供对硬件的显式控制（内存空间、向量/网格布局、同步）。工程价值：pallas_call 允许在 JAX 的高层框架内嵌入底层 kernel，实现对 TPU 的精细控制（如显式分配 SMEM/VMEM、控制 tile 布局），从而在框架无法自动优化的场景获得极致性能；同时保持与 JAX 生态的集成（数据以 JAX 数组传递、可编译为 XLA）。它为"需要底层控制但不想脱离 JAX"的场景提供了通道，让研究者能对 TPU 做深度调优。工程价值在于弥合"高层框架便利"与"底层硬件控制"之间的鸿沟。

pallas_call 的价值是"在 JAX 内写 TPU 底层 kernel"。它提供显式硬件控制，同时保持框架集成，适合需要深度调优 TPU 的场景。

• 理解 Mosaic 编译器
• 掌握从 Pallas 到 XLA HLO 的生成
• 理解工程价值

Mosaic 是 Pallas 的底层编译器，把 Pallas 描述的 TPU kernel 编译为 XLA HLO（或更低层的 TPU 指令），从而利用 TPU 的硬件特性。工程价值：Pallas 提供高层 kernel 描述，Mosaic 负责把 Pallas 的运算、内存布局与调度翻译为 XLA HLO，使 Pallas kernel 能进入 XLA 的编译与优化流程（如融合、布局、代码生成），并获得 TPU 的高效执行。Mosaic 让 Pallas 的"低层编程"与 XLA 的"编译优化"结合：用户用 Pallas 控制关键 Kernel，Mosaic 生成 HLO 并交给 XLA 优化，兼顾控制力与性能。工程价值在于打通 Pallas 抽象与 XLA 编译栈，让 TPU 深度定制可复用 XLA 的成熟优化。

Mosaic 是 Pallas 的编译器，把 Pallas kernel 生成 XLA HLO。它的价值是让 Pallas 的低层控制与 XLA 的编译优化结合，兼顾控制力与性能。

• 理解 TPU 内存层次
• 掌握显式分配的机制
• 理解工程价值

Pallas 允许用户显式分配 TPU 的多层内存空间：SMEM（片上标量内存，小且快）、VMEM（片上向量内存，更大）、HBM（片外高带宽内存，容量最大）。通过显式分配，程序员可精确定义数据放置在哪个内存层级，控制数据在 SMEM/VMEM/HBM 间的流动。工程价值：显式分配让 kernel 能针对 TPU 的片上内存做定制优化（如把热数据放 VMEM 提升带宽、用 SMEM 做标量/同步），避免编译器默认布局的低效；同时提供对内存带宽与容量的精细调度，实现接近硬件上限的性能。它为"需要精确控制内存布局的 TPU kernel"提供表达能力，是 Pallas 深度优化 TPU 的关键能力。

显式内存分配是 Pallas 深度调优 TPU 的核心能力。理解 SMEM/VMEM/HBM 的层次与分配，就能针对 TPU 内存带宽做精细优化。

• 理解 TMA 机制
• 掌握 hardware descriptor
• 理解工程价值

TMA（Tensor Memory Accelerator）是 Hopper/Blackwell 引入的硬件单元，用于加速全局内存与共享内存之间的张量数据搬运，通过 hardware descriptor（硬件描述符）描述要搬运的张量（形状、布局、stride、目标地址等），由硬件异步执行传输。工程价值：TMA 把数据搬运卸载给专门硬件，释放 SM 的计算资源，让数据加载与计算并行（流水线重叠）；硬件描述符支持复杂的多维张量布局与 swizzle，减少其他线程参与的搬运开销；配合 cp.async.bulk 与 mbarrier 实现异步批量传输，大幅提升数据吞吐。TMA 让 kernel 能用"描述一次、硬件搬运"的方式高效加载数据，是高性能 GEMM/Attention 的关键，尤其配合 WGMMA 与 pipeline 流水线。

TMA 的价值是"硬件卸载数据搬运 + 描述符驱动的异步传输"。它释放计算资源、支持复杂布局与流水线，是 Hopper/Blackwell 高性能算子的关键。

• 理解 cutlass-python
• 掌握 Python DSL 编写 kernel
• 理解工程价值

cutlass 3.x 提供 cutlass-python，这是一个 Python DSL，让用户用 Python 描述 GEMM 等 kernel 的配置（tile 大小、warp 布局、指令选择、内存流水），并生成对应的 cutlass C++ kernel 代码。工程价值：cutlass-python 把 cutlass 的底层 C++ 模板能力用 Python 表达，降低编写高性能 kernel 的门槛，让不熟悉 C++ 模板的开发者也能配置 cutlass 的高性能组件；同时保留 cutlass 的性能优势（TMA、WGMMA、多层流水线），实现"Python 描述 + C++ 高性能实现"。它让 kernel authoring 从繁琐的 C++ 模板编程变为声明式 Python 配置，同时生成可编译的最终代码，兼顾易用性与性能。

cutlass-python 的价值是"用 Python 描述、生成 C++ 高性能 kernel"。它降低门槛但保留 cutlass 的性能能力，是 kernel authoring 的易用化路径。

• 理解 Driver API 与 Runtime API 的差异
• 掌握动态加载的机制
• 理解工程价值

CUDA Driver API 提供比 Runtime API 更底层的控制。cuDeviceGet 获取设备句柄、cuModuleLoad 从 PTX/cubin 加载模块、cuLaunchKernel 直接启动 kernel。工程价值：Driver API 支持 kernel 动态加载（运行时从 PTX/cubin 加载），实现 JIT 编译、按需加载多样性 kernel、以及更精细的设备/上下文/流管理；相比 Runtime API 的静态/隐式加载，Driver API 让应用能灵活地把 kernel 代码作为资源动态加载与替换，适合需要动态扩展、插件化或热更新的场景。它还支持更细粒度的错误处理与资源控制。工程价值在于提供"运行时动态加载与底层控制"的能力，适合需要灵活内核管理与深度定制的应用。

Driver API 的价值是"底层控制 + 动态加载"。cuModuleLoad 动态加载 kernel、cuLaunchKernel 直接启动，支持 JIT 与插件化，比 Runtime API 更灵活。

• 理解 Pallas 与 Triton 的定位
• 掌握 CUDA kernel 编写能力
• 理解工程边界

Triton 与 Pallas 都提供高层 kernel 编写，但定位不同。Triton 以块级抽象和跨 GPU（CUDA/ROCm）可移植为目标，主要面向 NVIDIA GPU（CUDA），通过 tl.load/store 与块级算子编写 kernel，编译器自动处理底层；Pallas 主要面向 TPU（也支持部分 GPU），通过 pallas_call 提供显式内存空间与更细粒度控制。在 CUDA kernel authoring 上，Triton 是主流选择（PyTorch 的 torch.compile 默认用 Triton 生成 CUDA kernel），生态成熟、集成好；Pallas 在 CUDA 上支持相对有限，主要强项在 TPU。工程边界：需要 CUDA 高性能 kernel 且追求跨平台时用 Triton；需要 TPU 深度控制或已在 JAX 生态时用 Pallas。二者边界是"硬件重点（CUDA vs TPU）与抽象层次（自动 vs 显式）"。

Triton 是 CUDA 生态的主流高层 kernel 工具，Pallas 强在 TPU。工程边界取决于目标硬件与所需控制粒度，理解二者定位即可选型。

• 理解 TMA descriptor 与 cluster
• 掌握跨 SM 异步加载
• 理解工程价值

TMA 支持通过 hardware descriptor 把数据异步加载到显式地址，包括 cluster shared memory（集群共享内存）。在 Hopper 的 thread block cluster 机制下，多个 SM 组成 cluster，TMA 可以跨 SM 把数据加载到 cluster 内其他 SM 的共享内存，实现跨 SM 的数据共享与协作。工程价值：TMA descriptor 驱动的跨 SM 异步加载让数据搬运由硬件异步完成，释放计算资源，并支持 cluster 内 SM 间高效交换数据（如 Attention 中共享 Q/K/V 块），减少软件同步与冗余搬运；配合 mbarrier 做同步，实现大规模协作 kernel 的高效流水线。跨 SM 异步 load 的价值在于让 cluster 协作的 kernel 能高效复用数据，是分布式/集群 kernel 优化的关键。

TMA descriptor 的价值是"硬件异步 + 跨 SM 共享内存加载"。它支持 cluster 内 SM 协作、减少同步开销，是集群级 kernel 高性能的关键。

• 理解 TMA、LDGSTS、cp.async 的差异
• 掌握异步批量复制机制
• 理解工程边界

cp.async（及 LDGSTS，即 LDG + STS 的异步拷贝指令）是较老的一代异步批量拷贝指令，把全局内存数据异步搬运到共享内存，由单个线程或少量线程发起；TMA 是新一代硬件单元，用描述符驱动大批量张量传输，卸载更彻底、支持多维布局与 swizzle。协同工程边界：cp.async/LDGSTS 适合小规模、需要细粒度控制的异步拷贝，兼容性好、可用于早期架构；TMA 适合大规模、复杂布局、需要高吞吐与流水线重叠的场景，Hopper/Blackwell 上更高效。实践中常结合：TMA 负责大块、规则的批量加载，cp.async 处理小块或特殊访问，二者与 mbarrier 配合实现以计算重叠的流水线。边界取决于数据规模、布局复杂度与硬件代数，TMA 是新一代优选，cp.async 用于兼容与细粒度。

cp.async/LDGSTS 是旧一代、细粒度异步拷贝，TMA 是新一代、描述符驱动的大批量传输。协同边界在规模、布局与硬件代数，TMA 是主推方向。

• 理解 cutlass 3.x 的多层抽象
• 掌握 tile/warp/instruction 级 codegen
• 理解架构适配价值

cutlass 3.x 针对 Hopper（SM90）与 Blackwell（SM100）架构，提供 tile / warp / instruction 三级抽象与 codegen。tile level 决定分块与数据流（如分配 tile 到共享内存、warp 布局），warp level 决定 warp 如何分配输出与加载，instruction level 决定使用哪些指令（如 WGMMA、TMA、cp.async）。工程价值：cutlass 3.x 把这三层编码为可组合的模板，针对 SM90/SM100 的硬件特性（TMA、WGMMA、mbarrier、cluster）自动生成高效 kernel，让开发者无需手写底层即可获得针对新架构优化的实现；同时支持按需求定制某一层（如更换 warp 布局），兼顾性能与灵活性。它让新架构（Hopper/Blackwell）的硬件特性被快速、可复用地产出到高性能 kernel 中。

cutlass 3.x 的价值是"三级抽象 + 新架构适配 codegen"。它把 tile/warp/instruction 编码为模板，自动利用 SM90/SM100 的 TMA/WGMMA 等特性，兼顾性能与可定制。

• 理解硬件感知原语
• 掌握 TMA/WGMMA/mbarrier/cluster 的作用
• 理解工程价值

cutlass 3.x 提供丰富的 hardware-aware primitive（硬件感知原语），封装 Hopper/Blackwell 的底层硬件特性：TMA（异步张量搬运）、WGMMA（warpgroup 矩阵乘）、mbarrier（共享内存屏障，用于同步异步传输）、cluster shape（线程块集群拓扑）。工程价值：这些原语把硬件细节抽象为可组合、可复用的组件，让开发者能直接利用 TMA 卸载搬运、WGMMA 高效矩阵乘、mbarrier 做异步同步、cluster 做跨 SM 协作，而无需手写底层指令；原语是"硬件感知"的，能针对特定架构自动选择最优配置。工程价值在于让高性能 kernel 的开发基于"硬件原语 + 组合"而非"手写指令"，提升开发效率与性能上限，是 cutlass 在 Hopper/Blackwell 上高性能的核心。

cutlass 3.x 的价值是"硬件原语组合"。TMA/WGMMA/mbarrier/cluster 封装了硬件能力，让开发者高效利用新架构特性，兼顾性能与开发效率。

• 理解 cuMemAllocAsync 的异步内存池
• 掌握与流/内存池的协同
• 理解工程价值

cuMemAllocAsync / cuMemFree 是 CUDA 异步内存池 API，把内存分配/释放与执行流关联，通过内存池（memory pool）复用内存，避免每次分配都同步到驱动。工程价值：异步内存分配解耦了分配与计算，减少同步阻塞与分配开销；内存池复用已释放的内存，降低频繁分配/释放的碎片与驱动往返，提升内存效率与性能；按流关联内存池，让分配与流内计算正确序列化，避免竞争。CUDA 12.x 演进这些 API 支持更细粒度控制与跨流共享。工程价值在于把"内存管理"从同步开销中解放出来，与 cudaGraph 等配合，减少推理/训练中的内存延迟尖峰。

cuMemAllocAsync 的价值是"异步 + 内存池复用"。它解耦分配与计算、复用内存、与流关联，降低分配开销与碎片，是高性能内存管理的关键。

• 理解流属性设置
• 掌握 priority 与内存访问策略
• 理解工程价值

cuStreamSetAttribute / cuStreamGetAttribute 允许设置/查询流的属性，包括 priority（流优先级）与 access policy window（内存访问策略窗口，如 persistence 的 L2 缓存命中范围）。工程价值：通过设置流优先级，可让关键 kernel（如关键路径计算）优先于低优先级任务执行，减少延迟波动；通过内存访问策略窗口，可把特定数据标记为 L2 持久缓存，提升热数据的缓存命中率，减少 HBM 访问。Driver API 提供这些细粒度控制，让应用能按需求定制流调度与缓存策略，优化延迟与吞吐。工程价值在于"让流与内存策略可编程"，支持对推理/训练关键路径做精细调优。

cuStreamSetAttribute 的价值是"流属性可编程"。priority 控制调度优先级，access policy 控制缓存策略，从而优化延迟与命中率。

• 理解 dynamic shape 处理
• 掌握 mark_dynamic 与 reduce-overhead
• 理解工程价值

torch.compile 处理 dynamic shape（动态形状）时，可通过 mark_dynamic 标记某维度为动态，避免编译器为静态假设重新编译；torch.compile(mode='reduce-overhead') 则通过 CUDA Graphs 把多个 kernel 捕获为图，减少 kernel launch 开销。工程价值：动态形状（如变长序列、动态 batch）在真实推理中常见，mark_dynamic 让编译器为动态形状生成通用代码，避免频繁 recompile 带来的性能抖动；reduce-overhead 模式利用 graph 捕获降低 GPU 上的 launch 开销，适合小 kernel 多的推理。二者结合，让 torch.compile 在动态输入与高 launch 开销场景下保持高效，是推理部署优化的实用手段。

dynamic shape 处理避免 recompile，reduce-overhead 用 graph 减少 launch 开销。两者针对推理的实际痛点，是 torch.compile 部署优化的关键配置。

• 理解 Inductor 生成 Triton kernel
• 掌握 GPU 协同机制
• 理解工程价值

TorchInductor 在 GPU 上通过 codegen 生成 Triton kernel：把 torch.compile 捕获的图优化后，映射为 Triton 块级代码，再由 Triton 编译为 CUDA。工程价值：Inductor 与 Triton 协同，让 PyTorch 模型能自动生成高性能 GPU kernel，无需手写；Inductor 处理算子融合、布局、内存访问与 autotune，Triton 负责块级 kernel 生成与底层编译，二者分工完成"图优化 → kernel 生成 → 硬件执行"。这种协同让 torch.compile 能高效利用 GPU，同时支持多后端（ROCm 等）。工程价值在于把 GPU kernel 的生成自动化、可移植化，是 PyTorch 性能提升的基石。

Inductor 与 Triton 协同实现"图优化 + kernel 生成"。Inductor 负责高层优化，Triton 负责块级生成，二者配合让 PyTorch 自动获得高效 GPU kernel。

• 理解 guard 机制
• 掌握 recompile 触发条件
• 理解工程边界

torch.compile 通过对输入图做 guard（守卫）来快速复用已编译 kernel：当输入形状、dtype、设备等与 guard 匹配时直接复用，否则触发 recompile（重新编译）。工程边界：guard failure 会触发重新编译，带来编译开销与延迟尖峰；recompile 频繁（如动态形状、不同 batch 频繁变化）会导致性能抖动。torch.compile 的 guard 覆盖有限，无法覆盖所有属性（如某些数据依赖），因此动态输入容易触发 recompile。工程实践：通过 mark_dynamic 预声明动态维、限制形状变化、或使用图缓存/模式限制来减少 recompile。理解 guard/recompile 的边界，是部署 torch.compile 时避免性能抖动的关键：静态形状受益大，高度动态输入需谨慎处理。

guard 决定复用 or recompile，动态输入触发 recompile 带来抖动。工程边界是"静态形状受益、动态输入需管理"，理解 guard 才能避免性能陷阱。

• 理解 cuGraph 的 capture/instantiate
• 掌握 Driver API 的协同
• 理解工程边界

cuGraph 通过图捕获（capture）把 stream 上的 kernel 依赖记录为图，再 instantiate（实例化）并可回放（launch），减少反复 launch 的 CPU 开销。CUDA Driver API（cuGraphExec 系列）提供更底层的图执行控制，如 cuGraphInstantiate、cuGraphExecKernelNodeSetParams 等。协同工程边界：capture 阶段用 stream 捕获 kernel 依赖，instantiate 阶段把图编译为可执行对象，回放阶段用 cuGraphLaunch 执行；Driver API 允许在 instantiate 后动态修改节点参数（如 cuGraphExecKernelNodeSetParams），无需重新 capture。边界在于：capture 要求工作在图捕获模式下、kernel 参数需在捕获时确定，动态修改需通过 Driver API 的节点 API 而非重新捕获。协同价值是"捕获一次、实例化、灵活回放"，减少 launch 开销并支持动态调度。

cuGraph 的 capture/instantiate 与 Driver API 的节点修改协同，实现"一次捕获、多次回放、动态调参"。边界是 capture 模式限制与参数修改需用底层 API。

• 理解距离向量
• 掌握方向向量
• 理解依赖刻画

在循环依赖分析中，若迭代 I 与迭代 J=I+d 之间存在数据依赖（I 写、J 读，或反之），则向量 d 称为距离向量（distance vector），刻画两次迭代在循环各维上的偏移。方向向量（direction vector）用符号（<, =, >）表示每个循环维上目标迭代相对源迭代的趋势，如 (≤,<) 表示第一维不大于、第二维严格小于。距离向量给出精确的偏移，方向向量给出包含可能性的符号抽象。二者用于判断依赖是否跨 tile 边界、能否并行：若距离向量在某个循环维上非零，则该维存在跨迭代依赖，不能直接并行；方向向量用于依赖的保守分析（如判断依赖方向是否违反并行化）。距离向量精确、方向向量保守，都是依赖分析刻画数据依赖的关键工具。

距离向量给出精确偏移，方向向量给出符号趋势。二者刻画依赖的"跨迭代关系"，用于判断并行性与变换合法性。

• 理解 loop fusion
• 掌握融合的依赖判断
• 理解合法性与收益

loop fusion 把多个可合并的循环融合为一个，减少循环开销并提升数据局部性（融合后数据可在同一循环内复用）。依赖分析判断融合是否合法：融合会改变循环的执行顺序，需确保不破坏原循环间的数据依赖。具体地，若循环 A 写的数据被循环 B 读，融合后要求：A 的所有迭代在 B 的对应迭代之前完成（对于前向依赖），或 B 读的旧值在 A 写前已读取（对于反向依赖）。依赖分析通过检查两个循环迭代间的依赖方向（距离向量）来判断：若存在反向依赖（A 写、B 读且 B 的迭代早于 A），融合后顺序改变会破坏依赖，则不合法；或需通过重排/部分融合来解决。若能保证依赖方向在融合后仍成立，则融合合法。合法融合提升局部性、减少访存，是常见的循环优化。

fusion 的合法性取决于依赖方向是否在融合后保持。独立循环（无依赖）最易融合，存在反向依赖则需谨慎。依赖分析保证了融合不破坏数据顺序。

• 理解 unroll-and-jam 的局部性改善
• 掌握 interchange 的局部性改善
• 理解与 tiling 的配合

loop unroll-and-jam 在展开外层循环的同时提内层循环（jam），让多个外层迭代的对应内层迭代在寄存器中合并执行，增加数据在寄存器中的复用（寄存器局部性），减少内存访问。loop interchange 交换循环顺序，把访问模式从"按行主序"调整到更好的缓存局部性（如把访问最校内层维的循环放到最内层），提升空间/时间局部性。与 tiling 的配合：interchange 先调整循环顺序让访存更连续，tiling 再分块提升块内复用，unroll-and-jam 进一步在寄存器层复用——三者构成"顺序调整（interchange）→ 分块（tiling）→ 寄存器复用（unroll-and-jam）"的优化流水，共同最大化缓存与寄存器局部性。理解各自改善的局部性层次，才能正确组合它们。

interchange 改空间局部性、tiling 改块级缓存复用、unroll-and-jam 改寄存器复用。三者配合形成完整的局部性优化流水线。

• 理解算术强度的计算
• 掌握与 ridge point 的比较
• 理解判断方法

Roofline 模型用算术强度（arithmetic intensity，AI = 总 FLOPs / 总字节数）判断 loop nest 的瓶颈。首先计算该 loop nest 的算术强度（总浮点运算量除以总读写字节数），再与设备的 ridge point（= 算力峰值 / 带宽峰值）比较：若 AI 大于 ridge point，则 loop nest 是 compute-bound（计算受限），性能上限由算力决定；若 AI 小于 ridge point，则是 memory-bound（访存受限），性能上限由带宽决定。判断后可指导优化：访存受限时减少访存（提升复用、融合、压缩），计算受限时优化计算效率（提升 ILP/SIMD/并行）。Roofline 模型的价值在于用算术强度这一可计算指标，快速定位 loop nest 的瓶颈类型，从而决定优化方向。

Roofline 用算术强度与 ridge point 比较判断瓶颈。AI 高→计算受限，AI 低→访存受限。理解比较方法就能快速定位 loop nest 的优化方向。