[Docs]: Fact-check T.tile.fill API by zhangfei-machao · Pull Request #1634 · tile-ai/tilelang-ascend

zhangfei-machao · 2026-08-19T09:16:06Z

文件改动
文件类型改动说明
tilelang/language/ascend_tile.py 修改重写 fill 函数 docstring：补充后端指令说明（AscendC Duplicate / PTO TEXPANDS）、dtype 支持范围（AscendC 不支持 int8/uint8，PTO 通过 B82B16Trait 支持）、value 自动 Cast 语义、UB/shared scope 支持；按 Google style 重构（6 行 → 18 行含 Note）
testing/python/language/test_tilelang_ascend_language_elementwise.py 修改扩展 fill 测试套件：dtype 覆盖（7 dtype × AscendC + 9 dtype × PTO 含 int8/uint8）、1D shape、BufferRegion 切片、特殊值（负数/INT32_MAX）、shared 内存 fill；修复 assert_close_npu 处理 int8/uint8 比较；修复 run_test_fill int8/uint8 ref dtype 与填充值类型
docs/language/tile_fill.md 新增校验完成版 API 文档（功能说明 / 函数原型 / 参数规格 / 约束条件 / 示例代码），含 per-backend dtype 表
测试用例
2.1 约束测试（动态验证，真机 910B3 / dav_c220）
验证项测什么
AscendC dtype 全覆盖 float16/float32/bfloat16/int16/uint16/int32/uint32 × ascendc
PTO dtype 全覆盖同上 + int8/uint8 × pto
AscendC int8/uint8 int8/uint8 × ascendc 编译
PTO int8/uint8 int8/uint8 × pto 真机运行
bfloat16 PTO 支持 bfloat16 × pto 真机运行
value dtype 不一致 int32 buffer + float32 value × ascendc/pto
1D shape 1024×1 × float/float16/int32 × ascendc/pto
BufferRegion 切片 a_ub[0:32, 0:32] 子区域 fill × ascendc/pto
shared(L1) fill T.alloc_shared 直接 fill × ascendc/pto
特殊值 float16=-1.5, float=3.14159, int32=-100/INT32_MAX, int16=-32000
fill 与 clear 等价 fill(buf, 0) == clear(buf)
2.2 语言测试 test_tilelang_ascend_language_elementwise.py（40 用例）
测试函数用例数测什么
test_fill 14 7 dtype × ascendc/pto（不含 int8/uint8）
test_fill_pto_int8 2 int8/uint8 × pto only
test_fill_1d 6 1D shape × float/float16/int32 × ascendc/pto
test_fill_buffer_region 4 BufferRegion 子区域 fill × float/float16 × ascendc/pto
test_fill_special_values 10 负数/大值/π × 5 dtype-value 组合 × ascendc/pto
test_fill_shared 4 shared(L1) 直接 fill × float/float16 × ascendc/pto
测试结果

约束测试：动态验证完成（真机 910B3 / dav_c220，见 2.1 表）
语言测试：40 PASSED（真机 910B3，298s）
回归测试：test_clear 4 PASSED + test_axpy/test_adds/test_subs 14 PASSED
语法检查：py_compile 通过
whitespace 检查：git diff --check 无错误

校验过程中发现的问题
"bfloat16 仅 AscendC 后端支持"错误：原文档后端差异说明称 bfloat16 仅 AscendC 支持。真机验证 PTO bfloat16 同样通过（TEXPANDS 原生支持 bfloat16_t）→ 2.3.2 改为 per-backend dtype 表，bfloat16 两后端均标"支持"。
"value dtype 需与 buffer 一致"约束不成立：原文档约束 1 称 value dtype 需与 buffer 一致。elem.cc Fill 构造函数在 dtype 不匹配时执行 Cast(dst->dtype, args[1])，真机 int32 buffer + float32 value 两后端均通过 → 约束 1 改为"value dtype 不需要一致，前端自动 Cast"。
AscendC int8/uint8 不支持但原文档 A2/A3 行列出：原文档 DataType 表 A2/A3 行将 int8/uint8 与其他 dtype 并列，仅在注释中说明"仅 PTO"。dav_c220/kernel_operator_vec_duplicate_impl.h:23 的 static_assert 不含 int8/uint8，编译直接失败 → 改为 per-backend 表，AscendC 标"不支持"。
PTO 忽略 Python 层推断的 size：Python fill 将 math.prod(buffer.shape) 作为 size 传入 intrinsic，AscendC codegen 将其作为 count 传给 Duplicate；PTO codegen 完全忽略该参数，改用 GetSliceInfo 的 valid_row/valid_col 推断填充范围 → 2.3.3 注明两后端 size 推断路径不同。
fill 支持 shared(L1) 内存：原文档约束 4 称"fill 不区分 UB 和 L1/L0 级别"，约束 5 称"仅支持片上 buffer"。真机 T.alloc_shared 直接 fill 两后端均通过 → 约束 4 明确支持 UB 和 shared；fragment(L0C) 未完整验证（T.copy(a_frag, a_ub) 路径失败，非 fill 本身问题）。
"32 字节对齐"约束保留：原文档约束 2 称 buffer 地址需 32 字节对齐。AscendC Duplicate 底层 vector_dup 以 32B 为 block unit，与约束一致 → 保留。
A5 int64/uint64 无法验证：原文档 A5 行列出 int64/uint64。当前环境仅 910B3（dav_c220），CANN 安装中无 dav_m510 头文件 → 文档注明"当前环境无法验证"，static_assert 不含 int64/uint64，预计 AscendC 不支持。
FACTCHECK_WORKFLOW.md 不存在：任务要求遵循该文档，但仓库中未找到。沿用 T.tile.clear 事实校验先例（commit cf485f0）的文档结构和测试模式。
验证命令

fill 全量测试

python -m pytest testing/python/language/test_tilelang_ascend_language_elementwise.py -k "fill" -v

→ 40 passed, 490 deselected, 298s

clear 回归

python -m pytest testing/python/language/test_tilelang_ascend_language_elementwise.py -k "test_clear" -v

→ 4 passed, 526 deselected, 41s

其他回归

python -m pytest testing/python/language/test_tilelang_ascend_language_elementwise.py -k "test_axpy or test_adds or test_subs" -v

→ 14 passed, 516 deselected, 116s

语法检查

python -m py_compile tilelang/language/ascend_tile.py
python -m py_compile testing/python/language/test_tilelang_ascend_language_elementwise.py

- Add docs/language/tile_fill.md with verified API specification - Expand fill tests: dtype coverage (7 dtypes AscendC + 9 PTO incl int8/uint8), 1D shapes, BufferRegion slices, special values, shared memory - Update fill docstring with backend behavior and dtype support - Fix assert_close_npu to handle int8/uint8 comparison - Correct original doc errors: bfloat16 supported on PTO (not AscendC-only), value dtype auto-cast (not required to match), per-backend dtype table

github-actions · 2026-08-19T09:16:15Z

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run bash format.sh in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work!

🚀

liggest · 2026-08-24T03:19:48Z

+生成的底层指令取决于后端：
+
+- **Ascend C 后端**：`AscendC::Duplicate<T>(dst, value, count)`（通过 `tl::ascend::Fill<T>` 模板调用）
+- **PTO 后端**：`TEXPANDS(dst, value)`


参考其它接口文档，这里不用暴露后端对应的接口吧

liggest · 2026-08-24T03:20:46Z

+- **Ascend C 后端**：`AscendC::Duplicate<T>(dst, value, count)`（通过 `tl::ascend::Fill<T>` 模板调用）
+- **PTO 后端**：`TEXPANDS(dst, value)`
+
+`T.tile.clear(buf)` 内部委托给 `fill(buf, 0)` 实现。


T.tile.clear 也没必要在这里提及吧

liggest · 2026-08-24T03:26:24Z

+| buffer | 输入/输出 | 待填充的 buffer | Buffer / BufferRegion | 必填 |
+| value | 输入 | 填充的标量值 | Python 标量或 PrimExpr | 必填 |


liggest · 2026-08-24T03:31:20Z

+> **类型说明**：
+>
+> - **Buffer**：通过 `T.alloc_ub`、`T.alloc_shared` 等分配的片上缓冲区
+> - **BufferRegion**：Buffer 的切片，fill 将仅填充切片覆盖的区域
+> - **value**：可以是 Python 标量（如 `10.0`、`5`）或 PrimExpr（如 `T.cast(10, "float32")`）。当 value 的 dtype 与 buffer 不一致时，前端会自动将其 Cast 到 buffer 的 dtype


Suggested change

> **类型说明**：

>

> - **Buffer**：通过 `T.alloc_ub`、`T.alloc_shared` 等分配的片上缓冲区

> - **BufferRegion**：Buffer 的切片，fill 将仅填充切片覆盖的区域

> - **value**：可以是 Python 标量（如 `10.0`、`5`）或 PrimExpr（如 `T.cast(10, "float32")`）。当 value 的 dtype 与 buffer 不一致时，前端会自动将其 Cast 到 buffer 的 dtype

> **类型说明**：

> - **tensor**：通过 `T.alloc_ub`、`T.alloc_shared` 等分配的缓冲区（Buffer），或其切片（BufferRegion）

> - **scalar**：单个元素值，可以是 buffer 元素访问（BufferLoad）或 Python 标量/表达式（PrimExpr）

也是参考其它接口文档，按 tensor、scalar 的分类来解释，这里提供一个例子

liggest · 2026-08-24T03:36:29Z

+    The generated hardware instruction is ``AscendC::Duplicate`` on the
+    AscendC backend and ``TEXPANDS`` on the PTO backend.
+


同接口文档，不用在文档字符串里暴露后端对应的接口吧

liggest · 2026-08-24T04:04:05Z

+
+#### 2.3.1 Buffer Scope
+
+`fill` 支持 **UB（Unified Buffer）** 和 **shared（L1）** 内存。底层 `AscendC::Duplicate` 和 PTO `TEXPANDS` 均操作 `LocalTensor<T>`。


Duplicate 中有提到，dst 的存储位置为 UB，没有写支持 L1

如果要在 L1 上填充，看起来 Ascend C 中要用 Fill 接口，这应该是一个框架缺口

另外两个勘误：
shared 和 L1 不等价，shared 是 tilelang 里面 Buffer 的 scope，L1 是 NPU 内的存储单元，当前框架里应该是把 L1 映射成了 shared.l1、UB 映射成了 shared.ub
TEXPANDS 操作的也不是 Ascend C 的 LocalTensor<T>，而是 PTO 内部自己的 Tile

不过说到底这里应该也和上面一样，不用暴露后端对应的接口吧😶‍🌫️

liggest · 2026-08-24T04:50:09Z

+| int8 | 不支持 | 支持 |
+| uint8 | 不支持 | 支持 |
+
+> **说明**：Ascend C 后端 dtype 支持范围由 `AscendC::Duplicate` 的 `static_assert` 约束（dav_c220 平台），不支持 int8/uint8。PTO 后端使用 `TEXPANDS`，通过 `B82B16Trait` 将 B8 类型转换为 B16，支持 int8/uint8。


这些验证细节不用写在面向用户的文档里吧

liggest · 2026-08-24T04:51:56Z

+
+> **说明**：Ascend C 后端 dtype 支持范围由 `AscendC::Duplicate` 的 `static_assert` 约束（dav_c220 平台），不支持 int8/uint8。PTO 后端使用 `TEXPANDS`，通过 `B82B16Trait` 将 B8 类型转换为 B16，支持 int8/uint8。
+>
+> **A5 平台**：int64/uint64 在当前环境（A2/A3）无法验证。`AscendC::Duplicate` 的 dav_c220 `static_assert` 不包含 int64/uint64，Ascend C 后端预计不支持。


文档里可以先不涉及 A5 相关的内容

liggest · 2026-08-24T06:26:41Z

+- size 由 buffer shape 自动推断（BufferRegion 时取 region extent 的乘积，Buffer 时取 shape 的乘积）
+- 无需显式传入 count 参数；Ascend C 后端将推断的 size 作为 count 传入 `Duplicate`，PTO 后端根据 tile 的 valid_row/valid_col 自行推断填充范围


这些实现细节也不用写在面向用户的文档里吧

liggest · 2026-08-24T06:30:21Z

+1. value 的 dtype 不需要与 buffer 的 dtype 一致：当两者不同时，前端自动将 value Cast 到 buffer 的 dtype
+2. buffer 地址需 32 字节对齐（硬件约束）
+3. size 由 buffer shape 自动推断，无需显式传入 count 参数
+4. fill 支持 UB 和 shared 内存；fragment（L0C/寄存器）级别的 fill 行为未经完整验证


对 L1 的支持如上面的回复（ https://github.com/tile-ai/tilelang-ascend/pull/1634/changes#r3840624739 ）所述，应该存在缺口
L0C 目前应该是不支持，可以双后端再测试一下 / 去对应接口的文档确认一下

- Remove backend API names (Duplicate/TEXPANDS) from docs and docstring - Use tensor/scalar terminology in parameter table per convention - Remove T.tile.clear mention from fill doc - Remove implementation details (static_assert, B82B16Trait, count/valid_row) - Remove A5 platform content - Remove Buffer Scope section (L1 support is framework gap per reviewer) - Clarify only UB is verified; L0C/shared unverified - Remove clear equivalence example

liggest reviewed Aug 24, 2026

View reviewed changes

liggest mentioned this pull request Aug 25, 2026

[Docs]: Fact-check T.tile.clear API #1628

Open

Navigation Menu

Uh oh!

[Docs]: Fact-check T.tile.fill API - #1634

[Docs]: Fact-check T.tile.fill API #1634
zhangfei-machao wants to merge 2 commits into
tile-ai:ascendc_ptofrom
zhangfei-machao:factcheck/fill

zhangfei-machao commented Aug 19, 2026

Uh oh!

github-actions Bot commented Aug 19, 2026

Uh oh!

liggest Aug 24, 2026 •

edited

Loading

Uh oh!

liggest Aug 24, 2026

Uh oh!

liggest Aug 24, 2026

Uh oh!

liggest Aug 24, 2026

Uh oh!

liggest Aug 24, 2026 •

edited

Loading

Uh oh!

liggest Aug 24, 2026 •

edited

Loading

Uh oh!

liggest Aug 24, 2026

Uh oh!

liggest Aug 24, 2026

Uh oh!

liggest Aug 24, 2026

Uh oh!

liggest Aug 24, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

		\| buffer \| 输入/输出 \| 待填充的 buffer \| Buffer / BufferRegion \| 必填 \|
		\| value \| 输入 \| 填充的标量值 \| Python 标量或 PrimExpr \| 必填 \|

		The generated hardware instruction is ``AscendC::Duplicate`` on the
		AscendC backend and ``TEXPANDS`` on the PTO backend.


		#### 2.3.1 Buffer Scope

		`fill` 支持 UB（Unified Buffer）和 shared（L1）内存。底层 `AscendC::Duplicate` 和 PTO `TEXPANDS` 均操作 `LocalTensor<T>`。

		- size 由 buffer shape 自动推断（BufferRegion 时取 region extent 的乘积，Buffer 时取 shape 的乘积）
		- 无需显式传入 count 参数；Ascend C 后端将推断的 size 作为 count 传入 `Duplicate`，PTO 后端根据 tile 的 valid_row/valid_col 自行推断填充范围

Uh oh!

Conversation

zhangfei-machao commented Aug 19, 2026

fill 全量测试

→ 40 passed, 490 deselected, 298s

clear 回归

→ 4 passed, 526 deselected, 41s

其他回归

→ 14 passed, 516 deselected, 116s

语法检查

Uh oh!

github-actions Bot commented Aug 19, 2026

Uh oh!

liggest Aug 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

liggest Aug 24, 2026

Choose a reason for hiding this comment

Uh oh!

liggest Aug 24, 2026

Choose a reason for hiding this comment

Uh oh!

liggest Aug 24, 2026

Choose a reason for hiding this comment

Uh oh!

liggest Aug 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

liggest Aug 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

liggest Aug 24, 2026

Choose a reason for hiding this comment

Uh oh!

liggest Aug 24, 2026

Choose a reason for hiding this comment

Uh oh!

liggest Aug 24, 2026

Choose a reason for hiding this comment

Uh oh!

liggest Aug 24, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

liggest Aug 24, 2026 •

edited

Loading

liggest Aug 24, 2026 •

edited

Loading

liggest Aug 24, 2026 •

edited

Loading