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.