2026-09-30 的官方发布文章，列出了本次开源的六个组件与各自仓库地址。 本部分的所有章节都以这篇文章给出的清单为范围基线。

用 TileLang 写的生产算子库，同一份算法提供 CUDA 与 Ascend 两个后端。 是「一份算法、两个后端」写法的最佳教材，也是本教程第 33、34 章的主线。

DeepGEMM 的昇腾实现，支持 BF16/FP8/FP4 GEMM、MQA logits、MegaMoE 与 mHC prenorm。 主体是 Ascend C 与 MAD 原语抽象，mHC prenorm 由 TileLang 实现，适合讲两条路线的分工。

昇腾上的专家并行通信库：EP dispatch/combine、FP8 dispatch、延迟 epilogue， 以及 PP/Bucket/Engram 等通信原语。README 给出了 EP8–EP128 的实测带宽。

DeepGEMM-Ascend 的对照物。官方明确说 Ascend 版「完全兼容 DeepGEMM 的 API」， 所以两边对照着读才能看清哪些是平台无关的接口、哪些是平台相关的实现。

DeepEP-Ascend 公开 API 的对齐目标（V2.5 的 EPBuffer 接口）。

DeepSeek 的注意力算子库。2026-09-30 增加了昇腾 950 的稀疏 prefill/decode 算子， 并附一篇昇腾实现的技术报告（docs/20260930-ascend-prefill-deep-dive-zh.md）。

昇腾稀疏注意力算子的设计说明：1C2V 的推导、rescale 放哪、ND2NZ 时机、 gather2、高速 ND2NZ、icache 与 L1 bank 配平。本教程第 36 章的主要依据。

DSA 与采样器用的 TopK 算子，CUDA 与昇腾双平台。算法（阈值筛选 + RadixSelect） 与实现技巧分开讲，是「算法先行」的范例。

含随机块序下的期望复杂度证明，以及 CUDA 侧用 PTX 与位运算压低常数的技巧。

DeepGEMM-Ascend 与 DeepEP-Ascend 共用的运行时编译层。 和 TileLang 自己的 JIT 是同一层抽象的两个实现，值得对照。

TileKernels 里 Engram 算子的来源论文与参考实现。仓库内含 Engram_paper.pdf。

TileKernels 里 Engram 算子族的来源论文。给出 MoE（条件计算）与 Engram（条件记忆） 之间的容量分配 scaling law，并说明为什么确定性寻址可以把大表卸载到 host memory。 本地文本抽取见 14-deepseek-ascend-stack/raw/paper-deepseek-engram.md。

TileKernels 的 mhc/ 算子族与 DeepGEMM-Ascend 的 HC prenorm GEMM 的论文依据。 第 4.3 节（Kernel Fusion / Recomputing / Overlapping Communication in DualPipe） 正是这几个 kernel 与 DualPipe 的设计来源。本地文本抽取见 14-deepseek-ascend-stack/raw/paper-deepseek-mhc.md，PDF 为 paper-deepseek-mhc.pdf。

DeepSeek Sparse Attention（DSA）的首次公开实现。FlashMLA 的稀疏算子与 DeepSelect 的 TopK 都是为它服务的，先读这里再看算子更顺。

双向流水并行算法，用来讲「通信与计算重叠」在训练侧的组织方式。

专家并行负载均衡器。DeepEP-Ascend 的 lb_prefetch_weights / lb_reduce_grads 接口就是为它预留的。

昇腾 C/C++ 与 Ascend C 的 clangd 配置。写 Ascend C 侧代码时的补全与 lint 基础。

The grid/block/warp/thread model, the memory hierarchy, and the async primitives (cp.async, TMA, mbarrier). Everything TileLang's "expert level" interface exposes is defined here first.

Practical performance model: occupancy, coalescing, arithmetic intensity. The baseline against which kernel authors judge a tile configuration.

The instruction set TileLang ultimately emits. Read the sections on `mma`/`wgmma`/`tcgen05.mma`, `cp.async.bulk.tensor` (TMA), and the memory consistency model when you want to know what a TileLang intrinsic expands to.

Vendor tuning advice for sm90: TMA, clusters, distributed shared memory. The direct target of TileLang's cluster/TMA programming guide.

sm100/sm120 specifics: tensor memory (TMEM), tcgen05 MMA, block scaling. TileLang already has tcgen05 examples and macro generators; this is the document that explains the constraints they encode.

The C++ tile library TileLang competes with and partly mirrors (CuTe layouts, MMA atoms, epilogues). Reading it makes `T.gemm`'s layout/policy parameters much less magical.

Layout algebra (shape/stride/layout composition, tiling and partitioning of threads and values). TileLang's layout inference solves the same problem symbolically; this is the concrete formulation to compare against.

The vendor overview of Hopper: TMA (descriptor-based bulk async copy), thread block clusters and distributed shared memory, and the fourth-gen tensor core. This is the hardware story behind the `T.copy` -> TMA lowering and the cluster/TMA programming guide.

What changed for programs on Blackwell (sm100/sm120): new tensor-core and tensor-memory (TMEM) semantics, tcgen05 instructions, and the compatibility rules. Pair with the Blackwell Tuning Guide; the earlier blog-style architecture overview we tried to cite 404s, so this official guide is used instead.

How to measure a kernel: SOL sections, memory workload analysis, warp stalls. Needed for any serious tuning claim about a TileLang kernel.

Measured latencies/bandwidths for Hopper rather than vendor claims. The numbers you want when sanity-checking a pipeline stage count.

The original methodology paper for the Hopper/Blackwell dissections; useful when you need to measure an unfamiliar machine yourself.

The intermediate step of the same series; retained for completeness of the hardware-model reading path.

Reference CUDA implementations (including Hopper TMA/WGMMA samples). When a TileLang codegen path looks odd, compare the generated CUDA against the vendor's own sample.

Source for the C++ tile abstractions TileLang's CuTe/PTX intrinsics mirror, including the tcgen05 and block-scaled MMA generators.

The main repository (also cloned in full at 02-tilelang-core-repo/raw/tilelang). The README, docs/ tree, examples/ and tilelang/ + src/ trees are the ground truth for every claim in this corpus.

The low-bit operator library from the same MSRA group. Several ideas and contributors carried over into TileLang; reading it clarifies where the project's low-precision focus came from.

Huawei's own TileLang+Ascend training material: TileLang and MLIR basics, Pythonic programming, and a hand-written FlashAttention operator. The most direct official bridge between the two halves of this corpus.

The dedicated Ascend adapter: two backend routes (Ascend C & PTO, and AscendNPU IR), a programming guide, per-operator tutorials, and DeepSeek V4 kernels. This is the single most important source for the Huawei side.

The MLIR-based Ascend path. Where the TVM/TIR route and the MLIR route diverge is exactly the kind of question a would-be backend contributor has to answer, so both sources are needed.

Infrastructure for standing up new TileLang backends. Read together with 08-frontier-unmerged/raw/github-issue-2115 (multi-backend refactoring) to understand how a third chip vendor is expected to plug in.

The roadmap that made the Ascend/Hygon/MUSA adapters possible: how TileLang intends to decouple new backends from the CUDA-centric core. Essential for anyone planning to contribute a backend.

An open RFC with large blast radius: leaving the forked TVM for upstream TVM plus tvm-ffi. Explains a lot of in-flight churn in the tree.

Representative of the fine-grained control RFCs: exposing cache hints on plain global copies. Shows the level at which TileLang contributors argue.

Semantics-level RFC on reductions and buffer ownership; the kind of problem you meet as soon as you try to add a new reduction op or a new backend.

Memory-order semantics of atomics across codegen paths — a good test case for whether you actually understand TileLang's lowering.

The 0.2.0 plan; useful for dating features and for seeing which pieces were considered core versus experimental at that point.

The non-NVIDIA backend roadmap (Metal/WebGPU/SPIR-V). Read it for the pattern a new target follows, even though this corpus ignores those chips.

The CuTe DSL backend line. It matters because it defines the boundary between TileLang's own lowering and NVIDIA's newer Python tile DSL (cuTile/CuTe DSL), which is the most serious competitor on the NVIDIA side.

The PR that brought the Ascend 950 backend into mainline TileLang. The diff and review discussion are the most direct description of the Ascend/codegen surface that exists.

Adds the PTO codegen target for Ascend 950 — the second of the two Ascend backend routes, and the one that connects to PTO-ISA/pyPTO.

In-flight Ascend feature: persistent kernels plus the pass that derives them. An example of frontier work that has not landed, exactly what the user asked to keep an eye on.

In-flight: automatic SIMT/vector-fragment lowering for Ascend. Shows how the "tile -> SIMD/SIMT region" boundary is still being negotiated.

In-flight: LLM agents that write and tune Ascend TileLang operators. Signals where the project expects operator development to go next.

The maintainer discussion of Ascend C&PTO versus AscendNPU IR. Tl;dr-level guidance for choosing a route, straight from the people who wrote both.

Huawei's own account of training a large MoE on Ascend, including the software-stack constraints. Vendor source, and the least bad description of how the Ascend stack behaves at scale.

Benchmark/eval for NPU kernel generation on Ascend; a concrete yardstick for the "LLM writes Ascend kernels" workflows that TileLang@Ascend is part of.

Entry point to CANN: GE/Graph Engine, runtime, compiler (bisheng), operator libraries, HCCL. Establishes which layer a given Ascend problem belongs to.

The C/C++ kernel language TileLang@Ascend's PTO route lowers to. Reading it first makes the comparison "what does TileLang remove from Ascend C" concrete.

Huawei's LLVM-based compiler for Host-Device heterogeneous code. The backend TileLang's generated source is handed to.

Vendor overview of the operator development options available on Ascend (Ascend C, TBE, Triton, and now TileLang), useful for positioning TileLang inside Huawei's own story.

Map of Huawei's open-sourced CANN components (Ascend C, CATLASS, PTO, graph engine, ...). Use it as the authoritative list of where the source lives.

The PTO (Programmable Tensor Operations) ISA specification — the tile-level instruction set that TileLang@Ascend's PTO backend targets, and the subject of the 知乎 author's article series.

Huawei's CUTLASS analogue for Ascend C. The baseline TileLang@Ascend GEMM kernels are benchmarked against.

A Pythonic interface and JIT compiler for PTO-ISA — i.e. Huawei's own tile-DSL experiment next to TileLang. Comparing the two designs is one of the more informative exercises in this corpus.

Formal definition of the PTO ISA (normative language / ASL model). The precise semantics behind the Ascend tile instruction set.

Tile-operation tutorials and examples for PTO-based programming; the hands-on half of the PTO spec.

Hand-written PTO kernels. These are the performance ceiling that a TileLang kernel on Ascend has to be compared against.

Megakernel for Gated DeltaNet written with PTO-ISA. A concrete "what the low-level tile ISA can express" example on Ascend.

Python-first hardware construction toolkit with MLIR compilation. Included because the 知乎 author's article set treats it as the hardware-generation end of the same PTO story, and it shows where the tile abstraction is heading beyond kernel writing.

The multi-level IR philosophy. TileLang keeps its own TVM/TIR path but the Ascend backend has an MLIR-based sibling (tilelang-mlir-ascend), so this is the common language between the two approaches.

Best single orientation to the whole compiler landscape (graph IR, operator level, multi-level IR, backends). Use it as the first reading for anyone who has only used PyTorch and never seen a compiler pass pipeline.

The polyhedral-scheduling ancestor of today's tensor DSLs. Shows the alternative design point (automatic scheduling) that TileLang consciously trades away for predictability and controllability.

The original Triton paper (DOI 10.1145/3315508.3329973), mirrored locally as triton-mapl-2019.pdf with a page-by-page text extraction in triton-mapl-2019.md. This is the design point TileLang is measured against: Triton treats the program as blocks over a tensor with a compiler-chosen schedule; TileLang keeps explicit tiles, layouts and memory placement in the language. Read it for what "tiled intermediate language" originally meant, and note that the 2019 language is much smaller than today's Triton.

The Triton implementation, i.e. TileLang's most common comparison target. The original paper is "Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations" (MAPL 2019, DOI 10.1145/3315508.3329973); it is not mirrored here because the open copies we probed 404 and the ACM PDF is bot-blocked, so cite the DOI and use this repo plus the docs for detail. The decisive contrast with TileLang is that Triton hides the tile and lets the compiler pick the schedule, while TileLang keeps explicit tiles, layouts and memory placement in the language.

Current Triton language/API reference, for side-by-side reading with the TileLang language guide.

NVIDIA 官方的 Python tile DSL 文档。它是 TileLang 在 NVIDIA 侧最直接的对照物： 同样用 Python 写 tile 程序，但绑定 CUTLASS/CuTe 的抽象与 NVIDIA 自己的编译路径。 第 5 章对比与第 24 章 NVIDIA 后端都用它来界定「TileLang 自己的边界在哪里」。

cuTile 的开源实现（Python 包）。NVIDIA 新一代 tile 编程模型，面向 Blackwell 一代的硬件特性；与 TileLang 属于同一抽象层次的竞争路线。第 1、5、24 章引用它 说明「tile DSL」这一波的产业走向。

Upstream TVM/TIR documentation: IR structure, schedules, target/device abstraction. TileLang is a TVM dialect in practice, so this is its standard library documentation.

Dialect and pass infrastructure. Required reading for the AscendNPU IR route, which is MLIR-based rather than TVM-based.

The ancestor of the "separate algorithm from schedule" idea that TVM adopted and TileLang re-specialised for AI tiles.

The graph-capture/Inductor pipeline that a TileLang kernel has to fit into when it is used from PyTorch. Explains the boundary TileLang does *not* cover (graph-level fusion and autotuning).

How to see the kernels Inductor generates. The practical way to decide when torch.compile is enough and when a TileLang kernel has to be dropped in.

The single 知乎 article the user pointed at, and the best Chinese-language account of TileLang's Ascend port. Covers the NV-vs-950 architectural mismatch table, the new Ascend-specific intrinsics (T.alloc_l1/l0a/l0b/l0c, T.SimtVF/T.SimdVF, T.dual_copy), layout inference and ND->NZ fractal normalization, and how Z3 is used to solve the multi-pipeline scheduling and buffer-multiplicity constraints that Ascend C leaves to the human. Full text and images live in 09-zhihu-da-zhi-ruo-yu/raw/answers/.

Community lecture series on GPU programming that repeatedly covers TileLang alongside CUDA/Triton. Good source of talk recordings and slides links for the learning-path document.

Official course: environment setup, programming basics, the developer programming model, performance-tuning/debug tools, and deployment. Recorded in Chinese; the companion repo is tile-ai/tilelang-ascend. (Only the landing page is archived here; this corpus does not mirror video.)

Ten graded exercises that teach TileLang by making you write kernels. The best existing starting point for the "learn by doing" track of a course.

Language server for TileLang kernels. Minor for learning, but it documents the developer-experience side of the frontend.

端到端章节的基础工具文档：模型加载、tokenizer、Trainer/训练参数、混合精度、 checkpoint 保存与恢复。第 17 章「用 HF 资产训练 7B」直接依赖它。 （原地址 https://huggingface.co/docs/transformers/index 在本机不可达：huggingface.co 直连与走本地代理均 SSL EOF，改用中国镜像 hf-mirror.com 抓取；镜像内容与官网一致，引用时以原地址为准。）

数据集加载、流式读取、tokenize 与打包（packing）的官方说明；训练章节数据管线的 权威出处。 （原地址 https://huggingface.co/docs/datasets/index 在本机不可达：huggingface.co 直连与走本地代理均 SSL EOF，改用中国镜像 hf-mirror.com 抓取；镜像内容与官网一致，引用时以原地址为准。）

单机多卡/多机训练的抽象层，解释 device placement、混合精度与梯度同步在 HF 生态里 怎么表达。写训练章节时需要它来说明「哪些部分不归 TileLang 管」。 （原地址 https://huggingface.co/docs/accelerate/index 在本机不可达：huggingface.co 直连与走本地代理均 SSL EOF，改用中国镜像 hf-mirror.com 抓取；镜像内容与官网一致，引用时以原地址为准。）

参数高效微调的官方文档；用于说明「不做全量预训练时，TileLang 该优化哪几个算子」。 （原地址 https://huggingface.co/docs/peft/index 在本机不可达：huggingface.co 直连与走本地代理均 SSL EOF，改用中国镜像 hf-mirror.com 抓取；镜像内容与官网一致，引用时以原地址为准。）

7B 规模、权重与配置完全开放、许可证清晰的开源模型，作为端到端章节的训练/推理 对象。模型卡给出结构参数（层数、hidden、GQA 配置），是算子选择与显存估算的依据。 （原地址 https://huggingface.co/Qwen/Qwen2.5-7B 在本机不可达：huggingface.co 直连与走本地代理均 SSL EOF，改用中国镜像 hf-mirror.com 抓取；镜像内容与官网一致，引用时以原地址为准。）

公开的预训练语料样例，支持流式读取，适合教学规模的训练实验；训练章节用它构造 数据管线。 （原地址 https://huggingface.co/datasets/HuggingFaceFW/fineweb-edu 在本机不可达：huggingface.co 直连与走本地代理均 SSL EOF，改用中国镜像 hf-mirror.com 抓取；镜像内容与官网一致，引用时以原地址为准。）

小规模、可在单卡上跑完的语料，用于「先把训练循环跑通再放大」的教学路径。 （原地址 https://huggingface.co/datasets/roneneldan/TinyStories 在本机不可达：huggingface.co 直连与走本地代理均 SSL EOF，改用中国镜像 hf-mirror.com 抓取；镜像内容与官网一致，引用时以原地址为准。）

模型卡散文里没有的结构参数（hidden 3584 / intermediate 18944 / 28 层 / 28 头 / 4 个 KV 头 / vocab 152064 / rope_theta 1e6 / bf16）。第 17、18 章的显存与 KV cache 账本直接依赖这几个数字，所以单独存档（含原始 JSON 与摘要 Markdown）。

PyTorch 官方的 LLM 训练参考实现（并行策略、checkpoint、显存优化）。第 17 章用它 说明「哪些部分可以换成 TileLang kernel，哪些部分不要动」。

可读性最强的 LLM 训练/微调参考实现之一，适合作为教学骨架；与 HF 生态衔接良好。

最小可读的训练循环实现。教学价值在于「整条链路没有魔法」；用它讲清楚训练循环的 结构与可替换点。

推理服务章节的权威参考：PagedAttention、continuous batching、KV cache 管理与 自定义算子接入方式。第 18 章依赖它。

昇腾上的 vLLM 硬件插件。用于说明昇腾推理部署的实际形态，也是「TileLang 算子如何 进入服务框架」的对照物。

Defines MLA and the MoE structure that TileLang's flagship examples (deepseek_mla, fusedmoe, deepgemm) implement.

Context for why DeepSeek's kernels (and hence TileLang's DeepSeek examples) became an industry reference point; kept for the workload-level narrative.

High-performance LLM operator library built on TileLang. The natural reference for "what production TileLang kernels look like", including integration with PyTorch/vLLM.

Tile-based runtime for ultra-low-latency LLM inference. Shows what the kernels are deployed into, and the launch/scheduling constraints that come back as compiler requirements.

The project's own benchmark harness. Needed if you intend to make performance claims or detect regressions in a contribution.

Companion repo of the TileSight analytical model; used to reason about tile choices without brute-force autotuning.

The reference paper for everything else in this corpus. Introduces the tile abstraction, the three-level programming interface (beginner / developer / expert), layout inference, and the "composable" claim: high-level tile ops coexist with thread primitives in one kernel, lowered through TVM TIR.

The canonical tile-based fusion argument: the roofline/IO analysis that gives TileLang's attention examples their reason to exist.

Where the block partitioning, warp partitioning and shared-memory layout choices in TileLang's flash attention examples come from.

The Hopper-specific design (ping-pong warp specialization, async TMA/MMA) that TileLang's Hopper examples and its warp-specialization pass target.

Same MSRA group that produced BitBLAS and drove TileLang's early low-precision kernels. Useful for the origin story: TileLang grew out of a line of work on low-bit operator generation, not out of general-purpose compiler research.

Tile-centric analytical model; the companion repo lives in the same tile-ai org. Use it when reasoning about *why* a tiling choice wins, instead of guess-and-measure tuning.

Empirical study of how DSL-generated kernels get evaluated in practice, including TileLang. Useful counterweight for the performance claims in vendor material.

Agentic kernel synthesis across platforms, i.e. the "LLM writes the TileLang" direction. Relevant because TileLang ships an agent-oriented toolchain (see tilelang/tools and the [Ascend] agent PRs in the main repo).

Sparse-attention kernel work in the DeepSeek/DSA lineage that is implemented with TileLang-style tile kernels; shows what "frontier operator" means here.

Sibling work to StreamIndex; documents the indexer/sparse-attention operator family that TileLang kernels target in 2026.

Reports its kernels as implemented in TileLang, including what had to be hand-written at the thread level. A concrete example of TileLang as the implementation vehicle for a research result.

DeepSeek's sparse-attention design; the operator family whose TileLang implementations (deepseek_nsa / dsa_* examples) are among the most demanding in-tree workloads.

The compiler TileLang is implemented on top of. Reading it explains the schedule/compute separation, the Halide-derived IR, and why TileLang can reuse TVM's lowering and target machinery instead of writing its own.

Defines the block/axis IR that TileLang emits into. Inductor, TVM 0.9+ and TileLang all speak this dialect, so it is the vocabulary in which TileLang passes (layout inference, pipeline transform, warp specialization) are written.

The search-based auto-scheduling line of TVM. TileLang's autotuner inherits its split/reorder/fuse vocabulary, so this is the right background for understanding what TileLang deliberately does *not* do (it keeps the user in control of the tiling decision).

The design behind TVM's MetaSchedule, which TileLang's autotuner sits on. Explains trace-based scheduling and the cost-model/search machinery that TileLang reuses for `autotune` runs on real hardware.