diff --git a/skills/ascendc/ascendc-translator/references/TileLang-AscendC-API-Mapping.md b/skills/ascendc/ascendc-translator/references/TileLang-AscendC-API-Mapping.md index 03f90595..3c69a9f2 100644 --- a/skills/ascendc/ascendc-translator/references/TileLang-AscendC-API-Mapping.md +++ b/skills/ascendc/ascendc-translator/references/TileLang-AscendC-API-Mapping.md @@ -5,6 +5,25 @@ 说明: - 对于部分 API,tensor-tensor 和 tensor-scalar 的 lowering 不同,分别列出。 +## Kernel 索引映射 + +| TileLang API / 变量 | AscendC API / 表达式 | 备注 | +| --- | --- | --- | +| `cid` | AIC 侧:`AscendC::GetBlockIdx()`;AIV 侧:`AscendC::GetBlockIdx() / AscendC::GetSubBlockNum()` | 物理 AI core id,通常范围为 `0..19`。AIC 侧 block id 已按物理 core 编号;AIV 侧同一个 physical core 下有多个 Vector sub-block,需要用 `GetSubBlockNum()` 折算出 core id。 | +| `vid` | `AscendC::GetSubBlockIdx()` | core 内 Vector-side lane id,通常范围为 `0..1`。用于在同一个 physical core 内做两路 Vector 分工。 | + +常见写法: + +```cpp +if ASCEND_IS_AIC { + coreIdx = AscendC::GetBlockIdx(); +} +if ASCEND_IS_AIV { + coreIdx = AscendC::GetBlockIdx() / AscendC::GetSubBlockNum(); + vid = AscendC::GetSubBlockIdx(); +} +``` + ## 数据搬运 | TileLang API | AscendC API | 备注 | diff --git a/skills/ascendc/tilelang-designer/references/BlockLevelDesign.md b/skills/ascendc/tilelang-designer/references/BlockLevelDesign.md index 0650b57f..1d4a397f 100644 --- a/skills/ascendc/tilelang-designer/references/BlockLevelDesign.md +++ b/skills/ascendc/tilelang-designer/references/BlockLevelDesign.md @@ -44,6 +44,22 @@ Block-level design 用来确定 kernel 的块级组织方式,不展开具体 2. 若这种切分明显损失并行度或数据访问效率,可引入可控的跨 block 归并。 3. 原始 layout 不利于分工时,可做轴合并、拆分或重排。 +### Vector 侧二级切分 + +当一个逻辑 block 被分配到某个物理 AI Core 后,还应继续考虑该 core 内 2 个 Vector core 如何分工。推荐在 block-level 阶段就把这种二级切分写清楚: + +```python +vec_num = 2 +sub_block_elems = block_elems // vec_num + +with T.Kernel(usedCoreNum, is_npu=True) as (cid, vid): + for localIdx in T.serial(tasksPerCore): + bx = cid * tasksPerCore + localIdx + elem_base = bx * block_elems + vid * sub_block_elems + # TODO(tile-level): + # - this Vector lane owns [elem_base, elem_base + sub_block_elems) +``` + ### 例子一:Matmul - 输出:`C[M, N]` diff --git a/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md b/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md index a394a152..dd0b68a3 100644 --- a/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md +++ b/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md @@ -78,10 +78,13 @@ with T.Kernel(block_num, is_npu=True) as (cid, vid): ... ``` -Typical usage: +On Ascend NPU, a common execution model is 20 physical AI cores. Each core has +two Vector cores and one Cube core. In TileLang Ascend kernels: -- `cid`: block or tile id -- `vid`: Vector-side split id when a kernel has Cube/Vector cooperation +- `cid` identifies the physical AI core, normally in the range `0..19`. +- `vid` identifies the Vector-side lane within that core, normally in the range `0..1`. +- For pure Vector kernels, use `cid` for coarse task partitioning across cores and + `vid` for the two-way Vector split inside each core. ### 2.4 Loops and Control Flow @@ -274,15 +277,13 @@ Scalar usage note for `T.tile.mul`/`T.tile.add`/`T.tile.sub`/`T.tile.div` and si - `src1` can be a scalar. - When `src1` is a scalar, its dtype must match `src0` dtype. -- Use explicit constructors such as `T.float32(src1)` or `T.bfloat16(src1)` before passing the scalar. +- Use explicit constructors such as `T.float32(src1)` before passing the scalar. Example: ```python alpha = T.float32(1.0 / 127.0) T.tile.mul(scale_ub, row_max_ub, alpha) -beta = T.bfloat16(0.5) -T.tile.mul(x_ub, x_ub, beta) ``` ### 6.2 Compare and Select