From d159bc9e891acdfa2e013f7f3005933cf9978bc2 Mon Sep 17 00:00:00 2001 From: wzz <1114140864@qq.com> Date: Thu, 30 Apr 2026 16:08:21 +0800 Subject: [PATCH 1/4] docs: refine TileLang scalar constructor guidance --- .../references/TileLangAscendProgrammingGuide.md | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md b/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md index a394a152..ae9ff91f 100644 --- a/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md +++ b/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md @@ -274,15 +274,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 From 582470e2f9c5fb0663a38bfe24e0fea1222ed12d Mon Sep 17 00:00:00 2001 From: wzz <1114140864@qq.com> Date: Thu, 30 Apr 2026 18:18:22 +0800 Subject: [PATCH 2/4] docs: clarify TileLang Ascend kernel indices --- .../TileLang-AscendC-API-Mapping.md | 20 +++++++++++++++++++ .../TileLangAscendProgrammingGuide.md | 9 ++++++--- 2 files changed, 26 insertions(+), 3 deletions(-) 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..fca31f32 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,26 @@ 说明： - 对于部分 API，tensor-tensor 和 tensor-scalar 的 lowering 不同，分别列出。 +## Kernel 索引映射 + +| TileLang API / 变量 | AscendC API / 表达式 | 备注 | +| --- | --- | --- | +| `with T.Kernel(block_num, is_npu=True) as (cid, vid)` | `KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2)` 或 `KERNEL_TYPE_MIX_AIC_1_1` + `AscendC::GetBlockIdx()` / `AscendC::GetSubBlockIdx()` | `block_num` 对应 launch 的 `blockDim`。常见 Ascend NPU 有 20 个物理 AI core，每个 core 有 2 个 Vector core 和 1 个 Cube core。 | +| `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/TileLangAscendProgrammingGuide.md b/skills/ascendc/tilelang-designer/references/TileLangAscendProgrammingGuide.md index ae9ff91f..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 From bf8bd84ab6ba9819796629c55ea86dc6b62f74dc Mon Sep 17 00:00:00 2001 From: wzz <1114140864@qq.com> Date: Thu, 30 Apr 2026 18:27:35 +0800 Subject: [PATCH 3/4] docs: simplify TileLang kernel index mapping --- .../references/TileLang-AscendC-API-Mapping.md | 1 - 1 file changed, 1 deletion(-) 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 fca31f32..3c69a9f2 100644 --- a/skills/ascendc/ascendc-translator/references/TileLang-AscendC-API-Mapping.md +++ b/skills/ascendc/ascendc-translator/references/TileLang-AscendC-API-Mapping.md @@ -9,7 +9,6 @@ | TileLang API / 变量 | AscendC API / 表达式 | 备注 | | --- | --- | --- | -| `with T.Kernel(block_num, is_npu=True) as (cid, vid)` | `KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_MIX_AIC_1_2)` 或 `KERNEL_TYPE_MIX_AIC_1_1` + `AscendC::GetBlockIdx()` / `AscendC::GetSubBlockIdx()` | `block_num` 对应 launch 的 `blockDim`。常见 Ascend NPU 有 20 个物理 AI core，每个 core 有 2 个 Vector core 和 1 个 Cube core。 | | `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 分工。 | From 49a48e5875a89d39140de5417af2cb08284282a4 Mon Sep 17 00:00:00 2001 From: wzz <1114140864@qq.com> Date: Thu, 30 Apr 2026 18:42:30 +0800 Subject: [PATCH 4/4] docs: add TileLang vector sub-block guidance --- .../references/BlockLevelDesign.md | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) 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]`