Scenario / 使用场景

使用 TileLang 开发 Ascend A2 上的 Chunk GDN/KDA 算子，需要确定各阶段如何划分并行任务、如何利用 Vector/Cube 完成三角矩阵求逆，以及如何将 KDA 的逐通道门控计算转化为矩阵乘法。

这些知识分散在示例源码中。缺少相关知识时，容易将存在跨 chunk 依赖的状态更新错误并行化，或使用低效的求逆方式、数值不稳定的门控分解。因此希望提炼为可检索、可复用的开发指南。

Domain / 所属领域

ops / 算子开发

Technology or scope / 技术或作用域

TileLang

Query attempt / 已尝试的检索

补充知识前已进行检索：

• 检索词：chunk
• 筛选条件：domain=ops、technology=tilelang、platform=a2、status=all
• 结果：未召回相关知识卡，TileLang 目录当时只有导航入口。

已有 AscendC Chunk GDN 知识不能直接替代 TileLang 指南，也未完整覆盖 KDA 的逐通道门控分解及其数值稳定性。

Expected knowledge / 期望补充的知识

希望补充三张简短指南：

1. 任务切分：说明哪些阶段可以按 chunk 并行，为什么状态更新需要按 head 划分任务、在任务内顺序遍历 chunk，并展示对应的 TileLang Kernel 网格。
2. 分块求逆：说明 32 阶小块逐行求解，以及 64/128 阶递归合并的计算方法和 Vector/Cube 分工。
3. KDA 门控分解：说明如何构造带门控的 Q/K 矩阵乘法输入，参考门控为何能够相消，以及参考值选择和中间量精度的考虑。

Candidate sources / 候选来源

Checklists / 提交确认

• 我已搜索现有知识和 Issue，确认当前需求尚未被覆盖。
• 我理解没有可靠来源或明确适用范围的内容可能不会进入知识库。