Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
Expand Up @@ -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 | 备注 |
Expand Down
16 changes: 16 additions & 0 deletions skills/ascendc/tilelang-designer/references/BlockLevelDesign.md
Original file line number Diff line number Diff line change
Expand Up @@ -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]`
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down Expand Up @@ -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
Expand Down