From 40b601fda31ce3d0cf014f262e74f615d34755a3 Mon Sep 17 00:00:00 2001
From: GeYu <1258170822@qq.com>
Date: Mon, 11 May 2026 21:12:11 +0800
Subject: [PATCH] Stabilize AscendC and TileLang skill guidance

Document TileLang import, scalar comparison, and allocation constraints, plus AscendC kernel entry and registration failure guidance. These notes make the generation flow more stable and help agents avoid spending time repeatedly debugging known issues.
---
 .../references/AscendCVerification.md         |  8 ++++++++
 .../references/dsl2Ascendc.md                 |  2 +-
 skills/ascendc/tilelang-designer/SKILL.md     | 19 +++++++++++++++++++
 .../TileLangAscendProgrammingGuide.md         |  3 +++
 4 files changed, 31 insertions(+), 1 deletion(-)

diff --git a/skills/ascendc/ascendc-translator/references/AscendCVerification.md b/skills/ascendc/ascendc-translator/references/AscendCVerification.md
index dc8aebc7..5f9c7021 100644
--- a/skills/ascendc/ascendc-translator/references/AscendCVerification.md
+++ b/skills/ascendc/ascendc-translator/references/AscendCVerification.md
@@ -64,6 +64,14 @@ import _matmul_leakyrelu_ext as _ext
 
 ---
 
+### 验证失败排查思路
+
+编译通过但运行时失败，常见根因：
+
+**Kernel 注册失败** — 若日志中出现 `RegisterAscendBinary` 相关报错，检查 `kernel/` 下是否存在"一文件多 entry"或符号命名冲突。
+
+---
+
 ## 已知 NPU 精度问题与 Workaround
 
 ### 1. `torch.cumsum` float16 2D tensor `dim=0` 非确定性 bug
diff --git a/skills/ascendc/ascendc-translator/references/dsl2Ascendc.md b/skills/ascendc/ascendc-translator/references/dsl2Ascendc.md
index bb4d3bd7..7362e288 100644
--- a/skills/ascendc/ascendc-translator/references/dsl2Ascendc.md
+++ b/skills/ascendc/ascendc-translator/references/dsl2Ascendc.md
@@ -21,7 +21,7 @@
    提供调度、数据搬运、workspace 管理等通用能力。
 
 3. **Kernel 入口**：`xxx.cpp`
-   定义 `__global__ __aicore__` kernel 和 `extern "C"` launch 函数。
+   定义 `__global__ __aicore__` kernel 和 `extern "C"` launch 函数。每个 `.cpp` 文件只能有一个 kernel 入口；若需多个（如不同 dtype、tile 策略、任务类型），拆分为多个 `.cpp`，kernel 符号加统一前缀避免冲突。
 
 4. **主 Kernel 类与计算子模块**：一个或多个 `*.h`
    主 `Kernel` 类负责 `Init()` / `Process()` 主流程，管理 GM tensor、调度和流水。若 TileLang 中存在多个 `T.prim_func`，将对应的主 `Kernel` 类拆到多个独立头文件中，例如 `xxx_merge_n_kernel.h`、`xxx_single_row_kernel.h`。若算子属于 C/V 融合算子，或者 TileLang 设备侧存在多个有明确职责分工的 `Scope`，则可在这一部分下继续按计算阶段拆分子模块，例如 `matmul.h`、`leakyrelu.h`；通常每个 `Scope` 对应一个子模块，职责应与原 TileLang 设计中的计算阶段一一对应。对于纯 Vector 算子，或者虽然有 host / queue / buffer 管理但设备侧只有单个 Vector 计算阶段 / 单个 Vector `Scope` 的简单算子，主 `Kernel` 类本身通常就承载全部计算逻辑，不再额外拆分子模块。
diff --git a/skills/ascendc/tilelang-designer/SKILL.md b/skills/ascendc/tilelang-designer/SKILL.md
index 63c1064d..4dda2702 100644
--- a/skills/ascendc/tilelang-designer/SKILL.md
+++ b/skills/ascendc/tilelang-designer/SKILL.md
@@ -37,6 +37,25 @@ argument-hint: >
 └── <other_tasks>/        # 其他历史任务，可作为参考实现
 ```
 
+## model_new_tilelang.py 导入规范
+
+为保证 `model_new_tilelang.py` 在本地 AST 退化检测和远端运行时都能正确解析 `design.tile_level` 下的 kernel builder，必须同时满足以下两点：
+
+1. **在 `model_new_tilelang.py` 顶部添加 `sys.path.insert`**，将 `{output_dir}` 加入 Python 路径：
+   ```python
+   import sys
+   import os
+   sys.path.insert(0, os.path.dirname(os.path.abspath(__file__)))
+   ```
+   然后再写 `from design.tile_level.xxx import yyy`。
+
+2. **创建 `__init__.py` 使 design 成为 Python 包**。在生成 block_level / tile_level 设计的同时，必须创建以下空文件（或仅含注释）：
+   - `{output_dir}/design/__init__.py`
+   - `{output_dir}/design/tile_level/__init__.py`
+   - `{output_dir}/design/block_level/__init__.py`（如 block_level 下也有被导入的模块）
+
+这两个措施缺一不可：`sys.path.insert` 解决 Python 路径问题；`__init__.py` 解决包识别问题。
+
 ## Skill 参考资料
 本 skill 提供以下参考资料（位于 `@references/` 目录）：
 - `@references/BlockLevelDesign.md` — Block 层级设计指南
diff --git a/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md b/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md
index a394a152..a5cc4597 100644
--- a/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md
+++ b/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md
@@ -20,6 +20,7 @@ This guide focuses on that programming model and covers:
 ### 1.1 Programming Guidelines
 
 - Prefer `T.tile.*` APIs for compute whenever possible, and avoid scalar or element-by-element operations in hot paths.
+- **Do not use `T.min`, `T.max`, or similar tile-level APIs for scalar comparison.** These are tensor elementwise operations; for scalar min/max, use Python `min()`/`max()` (compile-time constants) or `T.if_then_else` (runtime variables).
 - Use `T.tile.broadcast` sparingly because it can consume large UB temporary space, and prefer row-wise or column-wise tile compute patterns when UB is constrained.
 - On the Vector side, in practice, you should copy inputs into UB and then cast them to `float32` at the beginning of `T.Scope("V")`, because this ensures better numerical stability and consistent behavior across the subsequent vector compute path.
 
@@ -101,6 +102,8 @@ Ordinary `if`, `else`, `while`, `break`, and `continue` can also be used in supp
 
 Ascend kernels in this guide only use L1, UB, and L0 storage.
 
+**Important**: The shape argument passed to any `T.alloc_*` API (`alloc_L1`, `alloc_ub`, `alloc_L0A`, `alloc_L0B`, `alloc_L0C`) must be a **compile-time integer constant expression**. Runtime variables are not allowed and will cause `Extent must be an integer constant` at compile time. If the actual data size varies at runtime, allocate the maximum fixed size and control the valid region via `T.copy` slices or conditional branches.
+
 ### 3.1 `T.alloc_L1`
 
 Allocates an L1 buffer.