已完成任务理解与源码调研。通读了 MindSpeed-Ops 的算子分层结构(api/ 接口层 + arch3x/ 实现层,经 is_arch35() 分平台分发)及 add、l2norm、rmsnorm 等典型 Triton 算子的实现范式,确认仓库中 TileLang 一列目前全为 ❌、且 api/tilelang、arch35/tilelang、tests/*/tilelang 骨架目录已预留。同时调研了 TileLang 昇腾后端(tilelang-ascend,lowering 到 AscendC 经 bisheng 编译,支持 A2/A3)的关键 API 映射与已知限制。据此选定 l2norm 前向作为首个打通流程的试点算子。

zeshengzong成员

7月2日关联了看板：@zeshengzong的看板 20260702

ascend-robot成员

7月3日关联了看板：MindStudio ISSUE管理

L liuzhexu成员

7月11日关联了里程碑：MindSpeed 26.2.0

kang

8月15日评论：

已提交pr

kang

8月15日评论：

实践文档

TileLang 算子迁移昇腾 NPU 实践文档（试点：l2norm 前向）

本文记录一次将 TileLang 算子迁移到昇腾 NPU 的完整实践，并沉淀为可复用的迁移 Skill。
试点算子：l2norm 前向。配套 Skill：tools/skills/tilelang-ascend-ops-migration/。

一、任务目标

将可加速开源仓库的 TileLang 源码迁移到昇腾 NPU（Atlas 800T A2/A3）上跑通并优化性能，产出一份「TileLang 算子迁移昇腾 NPU 的 Skill」，并以一个真实算子作为成功案例，完成精度/性能达标与 PR 合入。

二、结论速览

项	结果
试点算子	l2norm 前向（`x * rsqrt(sum(x²,-1)+eps)`）
环境	Atlas 800T A3 / openEuler 24.03 / CANN 9.0.0 / torch-npu 2.7.1 / tilelang-ascend 0.1.4
精度	双标杆 UT 28 项全过（fp16/bf16/fp32，L0 门限）
性能	大张量最高 5.26×（vs torch_npu eager，bf16 16384×4096）
交付	内核 + API + UT + ATK 用例 + 文档 + 迁移 Skill

三、迁移链路与仓库落地

TileLang 是面向 GPU 的 tile 级算子 DSL；tilelang-ascend 提供昇腾后端，将内核 lower 到 AscendC、经 bisheng 编译在 AICore 执行。迁移即把面向 GPU 的 TileLang 算子搬到昇腾 NPU 并达标。

落地到 MindSpeed-Ops 的分层结构：

层	文件	说明
接口层	`mindspeed_ops/api/tilelang/l2norm.py`	对外 API（l2norm/l2norm_fwd/L2Norm），JIT 内核按 (M,N,dtype) 缓存
实现层	`mindspeed_ops/arch35/tilelang/l2norm.py`	TileLang 内核（reduce→rsqrt→scale）
精度 UT	`tests/unit_tests/tilelang/test_l2norm.py`	双标杆，28 项
ATK 用例	`tests/atk_tests/tilelang/l2norm/`	yaml + generator + api
文档	`docs/tilelang/l2norm.md`	用法/精度/性能
Skill	`tools/skills/tilelang-ascend-ops-migration/`	迁移方法论 + 环境搭建

四、关键技术点

4.1 环境是最大门槛：glibc ≥ 2.38

tilelang-ascend 预编译 wheel 依赖 glibc ≥ 2.38。Ubuntu 22.04（glibc 2.35）镜像装上后 import 即报 GLIBC_2.38 not found；而源码编译需递归拉取 18 个 github 子模块（含定制 TVM，无国内镜像），实测不可行。改用 CANN + openEuler 24.03（glibc 2.38）镜像后解决。

完整依赖链（详见 Skill 的 QUICKSTART）：py3.11 venv → torch/torch_npu + PyYAML → tilelang-ascend → gcc-toolset-14（提供 GLIBCXX_3.4.32 的 libstdc++、g++、libgcc_s）→ python3-devel（JIT 编译 cython adapter 需 Python.h）→ TORCH_DEVICE_BACKEND_AUTOLOAD=0。

4.2 精度：硬件 rsqrt 需 Newton-Raphson 修正

昇腾硬件 T.tile.rsqrt 是快速低精度近似（相对误差约 3e-3），直接用无法通过仓库双标杆。对归约后的逆范数做 Newton-Raphson 迭代修正：

y_{n+1} = y_n * (1.5 - 0.5 * s * y_n²)

1 次迭代 → 相对误差约 1e-5（满足 fp16/bf16）
2 次迭代 → 约 2e-7（对齐 torch fp32）

迭代作用在已归约的 [ROWS,1] 向量上，开销极小。内核默认 2 次以覆盖 fp32。此结论对任意 TileLang 归一化类算子在昇腾上均适用。

4.3 精度验证：双标杆（dual-bench）

按仓库 docs/ops.md 规范，精度 UT 对齐双标杆：CPU float64 golden + torch 小算子参考，被测 TileLang 结果的误差比值需在门限内（MARE/MERE/RMSE 三指标，L0 档）。覆盖 fp16/bf16/fp32 × 多组 [T,D] × 多个 eps，共 28 项，全部通过。

五、性能数据

昇腾 910，TileLang 与 torch_npu eager 前向单次延迟对比：

shape	dtype	TileLang(µs)	torch_npu(µs)	加速比
[1024,512]	fp16	136.8	75.2	0.55×
[4096,1024]	fp16	135.5	80.4	0.59×
[16384,512]	bf16	138.4	112.9	0.82×
[16384,1536]	bf16	138.1	383.8	2.78×
[16384,4096]	bf16	331.7	1745.8	5.26×
[32768,2048]	bf16	357.0	1765.6	4.95×
[1024,1024]	fp32	138.7	62.9	0.45×

结论：TileLang 单次 kernel 有约 137µs 固定 launch 开销（批处理不摊薄，属设备侧启动成本）。小张量此开销占主导、不占优；大张量（线性注意力等场景的大激活张量）收益显著，最高约 5.26×。

六、复现步骤

申请 glibc ≥ 2.38 的昇腾环境（CANN + openEuler 24.03）。
按 QUICKSTART.md 搭建 tilelang 环境。
跑精度 UT：python -m pytest tests/unit_tests/tilelang/test_l2norm.py -q（期望 28 passed）。
跑性能对比脚本（见 Skill）。

七、已知限制与待补项

ATK 测试截图暂缺：仓库合入要求提供 ATK 精度/性能/内存截图（docs/ops.md 第 9 行）。ATK 安装包（atk*.whl）仅通过华为内部渠道分发，公开 gitcode 仓库（AscendTest/ATK）仅含使用文档、无源码/whl，当前无法获取安装包。ATK 用例已按 single_bm 标准编写完毕，待获取 whl 后补充截图。
竞品(GPU)精度对比暂缺：仓库要求 npu 算子能在竞品运行并提供与原 triton/TileLang 精度对比（docs/ops.md 第 10 行）。当前仅有昇腾 NPU 环境，无 GPU 环境，暂未提供。
仅前向：TileLang 后端当前仅实现前向；backward 抛 NotImplementedError，训练场景请用 triton 后端。

八、可复用产出：迁移 Skill

本次实践沉淀为 tilelang-ascend-ops-migration Skill：

SKILL.md：Phase 0–8 完整迁移工作流 + 真实 API 速查 + 常见坑速查表。
QUICKSTART.md：可复制粘贴的环境搭建命令。
README.md：Skill 说明与成功案例。

后续迁移其它 TileLang 算子（rmsnorm、softmax 等）可直接复用该 Skill，以 l2norm 四个产物文件为模板。

💻 需求背景、当前现状、期望实现的功能内容、具体的设计方案、以及测试方案

任务描述

验收标准

PR 合入

对接人

替代方案

补充说明