Star549
291
代码
代码
Issues59
Pull Requests118
Actions
项目讨论
Wiki
分析
Star549
291
已合并
revert: restore TMULA ISA API #1558
zhywang创建于 8月28日
revert: restore TMULA ISA API #1558
已合并
zhywang创建于 8月28日
从zhhywang/pto-isarename合入到cann/pto-isamaster
共 40 个文件变更+294-304
MOAT.xml+3-1
@@ -57,7 +57,9 @@
57 <!--filteritem type="filepath" name="projectroot/[a-zA-Z0-9]{20,}.sh" desc="Temp files"/-->57 <!--filteritem type="filepath" name="projectroot/[a-zA-Z0-9]{20,}.sh" desc="Temp files"/-->
58 </filefilter>58 </filefilter>
59 <filefilter name="copyrightPolicyFilter" desc="Filters for copyright header policies">59 <filefilter name="copyrightPolicyFilter" desc="Filters for copyright header policies">
60- <!--filteritem type="filename" name="*.uvwxyz" desc="Describe the reason for filtering scan results"/-->60+ <!-- JSON-compatible data files do not support comment headers. -->
61+ <filteritem type="filename" name="*.yaml" desc="skip yaml files"/>
62+ <filteritem type="filename" name="*.yml" desc="skip yml files"/>
61 <!--filteritem type="filepath" name="abcdefg/.*.uvwxyz" desc="Describe the reason for filtering scan results"/-->63 <!--filteritem type="filepath" name="abcdefg/.*.uvwxyz" desc="Describe the reason for filtering scan results"/-->
62 <!--filteritem type="filepath" name="projectroot/[a-zA-Z0-9]{20,}.sh" desc="Temp files"/-->64 <!--filteritem type="filepath" name="projectroot/[a-zA-Z0-9]{20,}.sh" desc="Temp files"/-->
63 </filefilter>65 </filefilter>
Mdocs/PTOISA.md+1-1
@@ -53,7 +53,7 @@ This page is the source-synchronized ISA index generated from `docs/isa/manifest
53| Elementwise (Tile-Tile) | [`TNEG`](isa/TNEG.md) | Elementwise negation of a tile. |53| Elementwise (Tile-Tile) | [`TNEG`](isa/TNEG.md) | Elementwise negation of a tile. |
54| Elementwise (Tile-Tile) | [`TREM`](isa/TREM.md) | Elementwise remainder of two tiles. |54| Elementwise (Tile-Tile) | [`TREM`](isa/TREM.md) | Elementwise remainder of two tiles. |
55| Elementwise (Tile-Tile) | [`TFMOD`](isa/TFMOD.md) | Elementwise fmod of two tiles. |55| Elementwise (Tile-Tile) | [`TFMOD`](isa/TFMOD.md) | Elementwise fmod of two tiles. |
56-| Elementwise (Tile-Tile) | [`TMULADDDST`](isa/TMULADDDST.md) | Elementwise ternary op: `src0 * src1 + dst`. |56+| Elementwise (Tile-Tile) | [`TMULA`](isa/TMULA.md) | Elementwise ternary op: `src0 * src1 + dst`. |
57| Tile-Scalar / Tile-Immediate | [`TEXPANDS`](isa/TEXPANDS.md) | Broadcast a scalar into a destination tile. |57| Tile-Scalar / Tile-Immediate | [`TEXPANDS`](isa/TEXPANDS.md) | Broadcast a scalar into a destination tile. |
58| Tile-Scalar / Tile-Immediate | [`TCMPS`](isa/TCMPS.md) | Compare a tile against a scalar and write per-element comparison results. |58| Tile-Scalar / Tile-Immediate | [`TCMPS`](isa/TCMPS.md) | Compare a tile against a scalar and write per-element comparison results. |
59| Tile-Scalar / Tile-Immediate | [`TSELS`](isa/TSELS.md) | Select between source tile and scalar using a mask tile (per-element selection for source tile). |59| Tile-Scalar / Tile-Immediate | [`TSELS`](isa/TSELS.md) | Select between source tile and scalar using a mask tile (per-element selection for source tile). |
Mdocs/PTOISA_zh.md+1-1
@@ -53,7 +53,7 @@
53| 逐元素（Tile-Tile） | [`TNEG`](isa/TNEG_zh.md) | Tile 的逐元素取负。 |53| 逐元素（Tile-Tile） | [`TNEG`](isa/TNEG_zh.md) | Tile 的逐元素取负。 |
54| 逐元素（Tile-Tile） | [`TREM`](isa/TREM_zh.md) | 两个 Tile 的逐元素余数，余数符号与除数相同。 |54| 逐元素（Tile-Tile） | [`TREM`](isa/TREM_zh.md) | 两个 Tile 的逐元素余数，余数符号与除数相同。 |
55| 逐元素（Tile-Tile） | [`TFMOD`](isa/TFMOD_zh.md) | 两个 Tile 的逐元素余数，余数符号与被除数相同。 |55| 逐元素（Tile-Tile） | [`TFMOD`](isa/TFMOD_zh.md) | 两个 Tile 的逐元素余数，余数符号与被除数相同。 |
56-| 逐元素（Tile-Tile） | [`TMULADDDST`](isa/TMULADDDST_zh.md) | 三元逐元素运算：`src0 * src1 + dst`。 |56+| 逐元素（Tile-Tile） | [`TMULA`](isa/TMULA_zh.md) | 三元逐元素运算：`src0 * src1 + dst`。 |
57| Tile-标量 / Tile-立即数 | [`TEXPANDS`](isa/TEXPANDS_zh.md) | 将标量广播到目标 Tile 中。 |57| Tile-标量 / Tile-立即数 | [`TEXPANDS`](isa/TEXPANDS_zh.md) | 将标量广播到目标 Tile 中。 |
58| Tile-标量 / Tile-立即数 | [`TCMPS`](isa/TCMPS_zh.md) | 将 Tile 与标量比较并写入逐元素比较结果。 |58| Tile-标量 / Tile-立即数 | [`TCMPS`](isa/TCMPS_zh.md) | 将 Tile 与标量比较并写入逐元素比较结果。 |
59| Tile-标量 / Tile-立即数 | [`TSELS`](isa/TSELS_zh.md) | 使用掩码 Tile 在源 Tile 和标量之间进行选择（源 Tile 逐元素选择）。 |59| Tile-标量 / Tile-立即数 | [`TSELS`](isa/TSELS_zh.md) | 使用掩码 Tile 在源 Tile 和标量之间进行选择（源 Tile 逐元素选择）。 |
Mdocs/coding/version-compatibility.md+1-0
@@ -147,6 +147,7 @@ Starting with PTO ISA v9.2.0, the public fused multiply-add instruction names us
147| Legacy name | Replacement | Semantics |147| Legacy name | Replacement | Semantics |
148| --- | --- | --- |148| --- | --- | --- |
149| `TFUSEDMULADD` | `TMADD` | `dst = src0 * dst + src1` |149| `TFUSEDMULADD` | `TMADD` | `dst = src0 * dst + src1` |
150+| `TMULADDDST` | `TMULA` | `dst = src0 * src1 + dst` |
150 151 
151Update C++ intrinsic calls, ISA references, generated diagrams, and ST testcase names to the replacement names. The legacy names are not retained as public wrappers after the rename.152Update C++ intrinsic calls, ISA references, generated diagrams, and ST testcase names to the replacement names. The legacy names are not retained as public wrappers after the rename.
152 153 
Mdocs/coding/version-compatibility_zh.md+1-0
@@ -147,6 +147,7 @@ Event 的详细模型说明请参考 [Event 编程模型](Event_zh.md)。
147| 历史名称 | 替代名称 | 语义 |147| 历史名称 | 替代名称 | 语义 |
148| --- | --- | --- |148| --- | --- | --- |
149| `TFUSEDMULADD` | `TMADD` | `dst = src0 * dst + src1` |149| `TFUSEDMULADD` | `TMADD` | `dst = src0 * dst + src1` |
150+| `TMULADDDST` | `TMULA` | `dst = src0 * src1 + dst` |
150 151 
151请同步更新 C++ intrinsic 调用、ISA 文档引用、生成图和 ST 用例名称。重命名后不再保留历史名称作为公开 wrapper。152请同步更新 C++ intrinsic 调用、ISA 文档引用、生成图和 ST 用例名称。重命名后不再保留历史名称作为公开 wrapper。
152 153 
Rdocs/figures/isa/TMULADDDST.svg→docs/figures/isa/TMULA.svg+2-2
@@ -1,4 +1,4 @@
1-<svg xmlns="http://www.w3.org/2000/svg" width="1200" height="720" viewBox="0 0 1200 720" role="img" aria-label="TMULADDDST tile operation diagram">1+<svg xmlns="http://www.w3.org/2000/svg" width="1200" height="720" viewBox="0 0 1200 720" role="img" aria-label="TMULA tile operation diagram">
2<defs>2<defs>
3 <marker id="arrow" markerWidth="12" markerHeight="12" refX="10" refY="6" orient="auto"><path d="M0,0 L0,12 L12,6 z" fill="#2D5BCE"/></marker>3 <marker id="arrow" markerWidth="12" markerHeight="12" refX="10" refY="6" orient="auto"><path d="M0,0 L0,12 L12,6 z" fill="#2D5BCE"/></marker>
4 <marker id="axisArrow" markerWidth="10" markerHeight="10" refX="8" refY="5" orient="auto"><path d="M0,0 L0,10 L10,5 z" fill="#64748b"/></marker>4 <marker id="axisArrow" markerWidth="10" markerHeight="10" refX="8" refY="5" orient="auto"><path d="M0,0 L0,10 L10,5 z" fill="#64748b"/></marker>
@@ -32,7 +32,7 @@ svg { font-family: Arial, Helvetica, sans-serif; }
32</style>32</style>
33<rect x="0" y="0" width="1200" height="720" class="frame" />33<rect x="0" y="0" width="1200" height="720" class="frame" />
34<rect x="24" y="24" width="1152" height="672" class="panel"/>34<rect x="24" y="24" width="1152" height="672" class="panel"/>
35-<text x="40" y="46" class="title">TMULADDDST</text>35+<text x="40" y="46" class="title">TMULA</text>
36<text x="40" y="72" class="subtitle">Elementwise ternary op: `src0 * src1 + dst`.</text>36<text x="40" y="72" class="subtitle">Elementwise ternary op: `src0 * src1 + dst`.</text>
37<text x="40" y="92" class="meta">Template: elementwise</text>37<text x="40" y="92" class="meta">Template: elementwise</text>
38<text x="1160" y="92" class="meta" text-anchor="end">Legend: outline=example; dashed=valid rows/cols (Rv,Cv); shaded=masked; r down / c right; ortho arrows=dataflow</text>38<text x="1160" y="92" class="meta" text-anchor="end">Legend: outline=example; dashed=valid rows/cols (Rv,Cv); shaded=masked; r down / c right; ortho arrows=dataflow</text>
Mdocs/isa/README.md+2-2
@@ -33,7 +33,7 @@ Migration guidance:
33- Replace `TSYNC(events...)` with ordinary event-based ordering: pass the event object to the consumer intrinsic, or call `WaitAllEvents(events...)` before the consumer when an explicit wait is required.33- Replace `TSYNC(events...)` with ordinary event-based ordering: pass the event object to the consumer intrinsic, or call `WaitAllEvents(events...)` before the consumer when an explicit wait is required.
34- `TSUBVIEW` is not a public ISA replacement. In-repository implementations may use `pto::detail::PtoSubTileView` as an internal helper; external code should express the data view through supported tile construction and public data movement APIs.34- `TSUBVIEW` is not a public ISA replacement. In-repository implementations may use `pto::detail::PtoSubTileView` as an internal helper; external code should express the data view through supported tile construction and public data movement APIs.
35- Replace ternary/scalar fused arithmetic forms with the corresponding primitive arithmetic sequence, such as `TADD`, `TSUB`, `TADDS`, `TSUBS`, `TMUL`, `TMADD`, and `TRELU`.35- Replace ternary/scalar fused arithmetic forms with the corresponding primitive arithmetic sequence, such as `TADD`, `TSUB`, `TADDS`, `TSUBS`, `TMUL`, `TMADD`, and `TRELU`.
36-- Rename legacy fused multiply-add APIs: `TFUSEDMULADD` to `TMADD`.36+- Rename legacy fused multiply-add APIs: `TFUSEDMULADD` to `TMADD`, and `TMULADDDST` to `TMULA`.
37- Replace fused add/ReLU/convert or add/dequant/ReLU forms with explicit arithmetic, conversion/dequantization, and `TRELU` steps.37- Replace fused add/ReLU/convert or add/dequant/ReLU forms with explicit arithmetic, conversion/dequantization, and `TRELU` steps.
38- Replace `TPairReduceSum` with the supported row/column reduction primitives that match the target layout.38- Replace `TPairReduceSum` with the supported row/column reduction primitives that match the target layout.
39- Do not call `TGET_SCALE_ADDR`; for AUTO-mode MX tests, bind the scale tile address from the data tile in test code before invoking the MX matmul primitive.39- Do not call `TGET_SCALE_ADDR`; for AUTO-mode MX tests, bind the scale tile address from the data tile in test code before invoking the MX matmul primitive.
@@ -78,7 +78,7 @@ Migration guidance:
78- [TNEG](TNEG.md) - Elementwise negation of a tile.78- [TNEG](TNEG.md) - Elementwise negation of a tile.
79- [TREM](TREM.md) - Elementwise remainder of two tiles.79- [TREM](TREM.md) - Elementwise remainder of two tiles.
80- [TFMOD](TFMOD.md) - Elementwise fmod of two tiles.80- [TFMOD](TFMOD.md) - Elementwise fmod of two tiles.
81-- [TMULADDDST](TMULADDDST.md) - Elementwise ternary op: `src0 * src1 + dst`.81+- [TMULA](TMULA.md) - Elementwise ternary op: `src0 * src1 + dst`.
82 82 
83## Tile-Scalar / Tile-Immediate83## Tile-Scalar / Tile-Immediate
84- [TEXPANDS](TEXPANDS.md) - Broadcast a scalar into a destination tile.84- [TEXPANDS](TEXPANDS.md) - Broadcast a scalar into a destination tile.
Mdocs/isa/README_zh.md+2-2
@@ -33,7 +33,7 @@
33- 将 `TSYNC(events...)` 替换为普通 event 顺序表达：把 event 对象传给消费端 intrinsic，或在确需显式等待时调用 `WaitAllEvents(events...)`。33- 将 `TSYNC(events...)` 替换为普通 event 顺序表达：把 event 对象传给消费端 intrinsic，或在确需显式等待时调用 `WaitAllEvents(events...)`。
34- `TSUBVIEW` 不提供公开 ISA 替代接口。仓内实现可使用 `pto::detail::PtoSubTileView` 作为内部 helper；外部代码应通过受支持的 Tile 构造和公开数据搬运 API 表达数据视图。34- `TSUBVIEW` 不提供公开 ISA 替代接口。仓内实现可使用 `pto::detail::PtoSubTileView` 作为内部 helper；外部代码应通过受支持的 Tile 构造和公开数据搬运 API 表达数据视图。
35- 将三元/标量融合算术形式替换为对应基础算术序列，例如 `TADD`、`TSUB`、`TADDS`、`TSUBS`、`TMUL`、`TMADD` 和 `TRELU`。35- 将三元/标量融合算术形式替换为对应基础算术序列，例如 `TADD`、`TSUB`、`TADDS`、`TSUBS`、`TMUL`、`TMADD` 和 `TRELU`。
36-- 将历史融合乘加接口重命名：`TFUSEDMULADD` 改为 `TMADD`。36+- 将历史融合乘加接口重命名：`TFUSEDMULADD` 改为 `TMADD`，`TMULADDDST` 改为 `TMULA`。
37- 将融合 add/ReLU/convert 或 add/dequant/ReLU 形式拆分为显式算术、转换/反量化和 `TRELU` 步骤。37- 将融合 add/ReLU/convert 或 add/dequant/ReLU 形式拆分为显式算术、转换/反量化和 `TRELU` 步骤。
38- 将 `TPairReduceSum` 替换为与目标 layout 匹配的现有行/列归约原语。38- 将 `TPairReduceSum` 替换为与目标 layout 匹配的现有行/列归约原语。
39- 不再调用 `TGET_SCALE_ADDR`；AUTO 模式 MX 测试可在调用 MX matmul 原语前，在测试代码中根据数据 Tile 绑定 scale Tile 地址。39- 不再调用 `TGET_SCALE_ADDR`；AUTO 模式 MX 测试可在调用 MX matmul 原语前，在测试代码中根据数据 Tile 绑定 scale Tile 地址。
@@ -78,7 +78,7 @@
78- [TNEG](TNEG_zh.md) - Tile 的逐元素取负。78- [TNEG](TNEG_zh.md) - Tile 的逐元素取负。
79- [TREM](TREM_zh.md) - 两个 Tile 的逐元素余数，余数符号与除数相同。79- [TREM](TREM_zh.md) - 两个 Tile 的逐元素余数，余数符号与除数相同。
80- [TFMOD](TFMOD_zh.md) - 两个 Tile 的逐元素余数，余数符号与被除数相同。80- [TFMOD](TFMOD_zh.md) - 两个 Tile 的逐元素余数，余数符号与被除数相同。
81-- [TMULADDDST](TMULADDDST_zh.md) - 三元逐元素运算：`src0 * src1 + dst`。81+- [TMULA](TMULA_zh.md) - 三元逐元素运算：`src0 * src1 + dst`。
82 82 
83## Tile-标量 / Tile-立即数83## Tile-标量 / Tile-立即数
84- [TEXPANDS](TEXPANDS_zh.md) - 将标量广播到目标 Tile 中。84- [TEXPANDS](TEXPANDS_zh.md) - 将标量广播到目标 Tile 中。
Rdocs/isa/TMULADDDST.md→docs/isa/TMULA.md+90-90
@@ -1,95 +1,95 @@
1-# TMULADDDST1+# TMULA
2- 2+ 
3-## Tile Operation Diagram3+## Tile Operation Diagram
4- 4+ 
5-![TMULADDDST tile operation](../figures/isa/TMULADDDST.svg)5+![TMULA tile operation](../figures/isa/TMULA.svg)
6- 6+ 
7-## Introduction7+## Introduction
8- 8+ 
9-Elementwise operation: `src0 * src1 + dst`.9+Elementwise operation: `src0 * src1 + dst`.
10- 10+ 
11-## Math Interpretation11+## Math Interpretation
12- 12+ 
13-For each element `(i, j)` in the valid region:13+For each element `(i, j)` in the valid region:
14- 14+ 
15-$$ \mathrm{dst}_{i,j} = \mathrm{src0}_{i,j} \* \mathrm{src1}_{i,j} + \mathrm{dst}_{i,j} $$15+$$ \mathrm{dst}_{i,j} = \mathrm{src0}_{i,j} \* \mathrm{src1}_{i,j} + \mathrm{dst}_{i,j} $$
16- 16+ 
17-## Assembly Syntax17+## Assembly Syntax
18- 18+ 
19-Synchronous form:19+Synchronous form:
20- 20+ 
21-```text21+```text
22-%dst = tmuladddst %src0, %src1 : !pto.tile<...>22+%dst = tmula %src0, %src1 : !pto.tile<...>
23-```23+```
24- 24+ 
25-### AS Level 1 (SSA)25+### AS Level 1 (SSA)
26- 26+ 
27-```text27+```text
28-%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>28+%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
29-```29+```
30- 30+ 
31-### AS Level 2 (DPS)31+### AS Level 2 (DPS)
32- 32+ 
33-```text33+```text
34-pto.tmuladddst ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)34+pto.tmula ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
35-```35+```
36-## C++ Intrinsic36+## C++ Intrinsic
37- 37+ 
38-Declared in `include/pto/common/pto_instr.hpp`:38+Declared in `include/pto/common/pto_instr.hpp`:
39- 39+ 
40-```cpp40+```cpp
41-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>41+template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>
42-PTO_INST RecordEvent TMULADDDST(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &...events);42+PTO_INST RecordEvent TMULA(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &...events);
43-```43+```
44- 44+ 
45-## Constraints45+## Constraints
46- 46+ 
47-- **Implementation checks**:47+- **Implementation checks**:
48- - `TileData::DType` must be one of: `float`, `half`.48+ - `TileData::DType` must be one of: `float`, `half`.
49- - Tile layout must be row-major (`TileData::isRowMajor`).49+ - Tile layout must be row-major (`TileData::isRowMajor`).
50- **Valid region**:50- **Valid region**:
51 - The op iterates over `dst.GetValidRow()` / `dst.GetValidCol()`.51 - The op iterates over `dst.GetValidRow()` / `dst.GetValidCol()`.
52 - CPU_SIM asserts that `dst`, `src0`, and `src1` have identical valid row and column counts.52 - CPU_SIM asserts that `dst`, `src0`, and `src1` have identical valid row and column counts.
53- In CPU_SIM, `half` multiplication is rounded to `half` before accumulation, and the accumulated result is rounded to53- In CPU_SIM, `half` multiplication is rounded to `half` before accumulation, and the accumulated result is rounded to
54 `half` again.54 `half` again.
55- 55+ 
56-## Examples56+## Examples
57- 57+ 
58-```cpp58+```cpp
59-#include <pto/pto-inst.hpp>59+#include <pto/pto-inst.hpp>
60- 60+ 
61-using namespace pto;61+using namespace pto;
62- 62+ 
63-void example() {63+void example() {
64- using TileT = Tile<TileType::Vec, float, 16, 16>;64+ using TileT = Tile<TileType::Vec, float, 16, 16>;
65- TileT a, b, out;65+ TileT a, b, out;
66- TMULADDDST(out, a, b);66+ TMULA(out, a, b);
67-}67+}
68-```68+```
69- 69+ 
70-## ASM Form Examples70+## ASM Form Examples
71- 71+ 
72-### Auto Mode72+### Auto Mode
73- 73+ 
74-```text74+```text
75-# Auto mode: compiler/runtime-managed placement and scheduling.75+# Auto mode: compiler/runtime-managed placement and scheduling.
76-%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>76+%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
77-```77+```
78- 78+ 
79-### Manual Mode79+### Manual Mode
80- 80+ 
81-```text81+```text
82-# Manual mode: resources must be bound explicitly before issuing the instruction.82+# Manual mode: resources must be bound explicitly before issuing the instruction.
83-# Optional for tile operands:83+# Optional for tile operands:
84-# pto.tassign %arg0, @tile(0x1000)84+# pto.tassign %arg0, @tile(0x1000)
85-# pto.tassign %arg1, @tile(0x2000)85+# pto.tassign %arg1, @tile(0x2000)
86-%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>86+%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
87-```87+```
88- 88+ 
89-### PTO Assembly Form89+### PTO Assembly Form
90- 90+ 
91-```text91+```text
92-%dst = tmuladddst %src0, %src1 : !pto.tile<...>92+%dst = tmula %src0, %src1 : !pto.tile<...>
93-# AS Level 2 (DPS)93+# AS Level 2 (DPS)
94-pto.tmuladddst ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)94+pto.tmula ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
95-```95+```
Rdocs/isa/TMULADDDST_zh.md→docs/isa/TMULA_zh.md+97-97
@@ -1,99 +1,99 @@
1-# TMULADDDST1+# TMULA
2- 2+ 
3-## 指令示意图3+## 指令示意图
4- 4+ 
5-![TMULADDDST tile operation](../figures/isa/TMULADDDST.svg)5+![TMULA tile operation](../figures/isa/TMULA.svg)
6- 6+ 
7-## 简介7+## 简介
8- 8+ 
9-三元逐元素运算：`src0 * src1 + dst`。9+三元逐元素运算：`src0 * src1 + dst`。
10- 10+ 
11-## 数学语义11+## 数学语义
12- 12+ 
13-对每个元素 `(i, j)` 在有效区域内：13+对每个元素 `(i, j)` 在有效区域内：
14- 14+ 
15-$$ \mathrm{dst}_{i,j} = \mathrm{src0}_{i,j} \* \mathrm{src1}_{i,j} + \mathrm{dst}_{i,j} $$15+$$ \mathrm{dst}_{i,j} = \mathrm{src0}_{i,j} \* \mathrm{src1}_{i,j} + \mathrm{dst}_{i,j} $$
16- 16+ 
17-## 汇编语法17+## 汇编语法
18- 18+ 
19-同步形式：19+同步形式：
20- 20+ 
21-```text21+```text
22-%dst = tmuladddst %src0, %src1 : !pto.tile<...>, !pto.tile<...>22+%dst = tmula %src0, %src1 : !pto.tile<...>, !pto.tile<...>
23-```23+```
24- 24+ 
25-### AS Level 1（SSA）25+### AS Level 1（SSA）
26- 26+ 
27-```text27+```text
28-%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>28+%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
29-```29+```
30- 30+ 
31-### AS Level 2（DPS）31+### AS Level 2（DPS）
32- 32+ 
33-```text33+```text
34-pto.tmuladddst ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)34+pto.tmula ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
35-```35+```
36- 36+ 
37-## C++ 内建接口37+## C++ 内建接口
38- 38+ 
39-声明于 `include/pto/common/pto_instr.hpp`：39+声明于 `include/pto/common/pto_instr.hpp`：
40-> 公共包含头为 `<pto/pto-inst.hpp>`，内部声明位于 `pto/common/pto_instr.hpp`。40+> 公共包含头为 `<pto/pto-inst.hpp>`，内部声明位于 `pto/common/pto_instr.hpp`。
41- 41+ 
42-```cpp42+```cpp
43-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>43+template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>
44-PTO_INST RecordEvent TMULADDDST(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &...events);44+PTO_INST RecordEvent TMULA(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &...events);
45-```45+```
46- 46+ 
47-## 约束47+## 约束
48- 48+ 
49-- **实现检查**:49+- **实现检查**:
50- - `TileData::DType` 必须是以下之一：`half`、`float`。50+ - `TileData::DType` 必须是以下之一：`half`、`float`。
51- - Tile 布局必须是行主序（`TileData::isRowMajor`）。51+ - Tile 布局必须是行主序（`TileData::isRowMajor`）。
52-- **通用约束**:52+- **通用约束**:
53- - Tile 位置必须是向量（`TileData::Loc == TileType::Vec`）。53+ - Tile 位置必须是向量（`TileData::Loc == TileType::Vec`）。
54- - 静态有效边界：`TileData::ValidRow <= TileData::Rows` 且 `TileData::ValidCol <= TileData::Cols`。54+ - 静态有效边界：`TileData::ValidRow <= TileData::Rows` 且 `TileData::ValidCol <= TileData::Cols`。
55- - 运行时：`dst`、`src0` 和 `src1` 的有效行列数必须相同。55+ - 运行时：`dst`、`src0` 和 `src1` 的有效行列数必须相同。
56- - 标量类型必须与 Tile 数据类型一致。56+ - 标量类型必须与 Tile 数据类型一致。
57- 该操作在 `dst.GetValidRow()` / `dst.GetValidCol()` 上迭代。57- 该操作在 `dst.GetValidRow()` / `dst.GetValidCol()` 上迭代。
58- CPU_SIM 的 `half` 路径先将乘积舍入为 `half`，再执行累加，并将累加结果再次舍入为 `half`。58- CPU_SIM 的 `half` 路径先将乘积舍入为 `half`，再执行累加，并将累加结果再次舍入为 `half`。
59- 59+ 
60-## 示例60+## 示例
61- 61+ 
62-```cpp62+```cpp
63-#include <pto/pto-inst.hpp>63+#include <pto/pto-inst.hpp>
64- 64+ 
65-using namespace pto;65+using namespace pto;
66- 66+ 
67-void example() {67+void example() {
68- using TileT = Tile<TileType::Vec, float, 16, 16>;68+ using TileT = Tile<TileType::Vec, float, 16, 16>;
69- TileT a, b, out;69+ TileT a, b, out;
70- TMULADDDST(out, a, b);70+ TMULA(out, a, b);
71-}71+}
72-```72+```
73- 73+ 
74-## 汇编示例（ASM）74+## 汇编示例（ASM）
75- 75+ 
76-### 自动模式76+### 自动模式
77- 77+ 
78-```text78+```text
79-# 自动模式：由编译器/运行时负责资源放置与调度。79+# 自动模式：由编译器/运行时负责资源放置与调度。
80-%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>80+%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
81-```81+```
82- 82+ 
83-### 手动模式83+### 手动模式
84- 84+ 
85-```text85+```text
86-# 手动模式：先显式绑定资源，再发射指令。86+# 手动模式：先显式绑定资源，再发射指令。
87-# 可选（当该指令包含 tile 操作数时）：87+# 可选（当该指令包含 tile 操作数时）：
88-# pto.tassign %arg0, @tile(0x1000)88+# pto.tassign %arg0, @tile(0x1000)
89-# pto.tassign %arg1, @tile(0x2000)89+# pto.tassign %arg1, @tile(0x2000)
90-%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>90+%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
91-```91+```
92- 92+ 
93-### PTO 汇编形式93+### PTO 汇编形式
94- 94+ 
95-```text95+```text
96-%dst = tmuladddst %src0, %src1 : !pto.tile<...>, !pto.tile<...>96+%dst = tmula %src0, %src1 : !pto.tile<...>, !pto.tile<...>
97-# AS Level 2 (DPS)97+# AS Level 2 (DPS)
98-pto.tmuladddst ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)98+pto.tmula ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
99-```99+```
Mdocs/isa/manifest.yaml+1-1
@@ -447,7 +447,7 @@
447 "notes": []447 "notes": []
448 },448 },
449 {449 {
450- "instruction": "TMULADDDST",450+ "instruction": "TMULA",
451 "category": "Elementwise (Tile-Tile)",451 "category": "Elementwise (Tile-Tile)",
452 "summary_en": "Elementwise ternary op: `src0 * src1 + dst`.",452 "summary_en": "Elementwise ternary op: `src0 * src1 + dst`.",
453 "summary_zh": "三元逐元素运算：`src0 * src1 + dst`。",453 "summary_zh": "三元逐元素运算：`src0 * src1 + dst`。",
Mdocs/mkdocs/gen_pages.py+1-1
@@ -302,7 +302,7 @@ NAV_PAGES_EN = [
302 "docs/isa/TLOG.md",302 "docs/isa/TLOG.md",
303 "docs/isa/TRECIP.md",303 "docs/isa/TRECIP.md",
304 "docs/isa/TPRELU.md",304 "docs/isa/TPRELU.md",
305- "docs/isa/TMULADDDST.md",305+ "docs/isa/TMULA.md",
306 "docs/isa/TCVT.md",306 "docs/isa/TCVT.md",
307 "docs/isa/TSEL.md",307 "docs/isa/TSEL.md",
308 "docs/isa/TRSQRT.md",308 "docs/isa/TRSQRT.md",
Mdocs/mkdocs/mkdocs.yml+1-1
@@ -97,7 +97,7 @@ nav:
97 - TLOG: docs/isa/TLOG.md97 - TLOG: docs/isa/TLOG.md
98 - TRECIP: docs/isa/TRECIP.md98 - TRECIP: docs/isa/TRECIP.md
99 - TPRELU: docs/isa/TPRELU.md99 - TPRELU: docs/isa/TPRELU.md
100- - TMULADDDST: docs/isa/TMULADDDST.md100+ - TMULA: docs/isa/TMULA.md
101 - TCVT: docs/isa/TCVT.md101 - TCVT: docs/isa/TCVT.md
102 - TSEL: docs/isa/TSEL.md102 - TSEL: docs/isa/TSEL.md
103 - TRSQRT: docs/isa/TRSQRT.md103 - TRSQRT: docs/isa/TRSQRT.md
Mdocs/mkdocs/src/manual/appendix-d-instruction-family-matrix.md+1-1
@@ -66,7 +66,7 @@ This appendix is generated from `docs/isa/manifest.yaml` and provides a source-s
66| Elementwise (Tile-Tile) | `TNEG` | `elementwise` | `dst, src0, src1` | `docs/isa/TNEG.md` |66| Elementwise (Tile-Tile) | `TNEG` | `elementwise` | `dst, src0, src1` | `docs/isa/TNEG.md` |
67| Elementwise (Tile-Tile) | `TREM` | `elementwise` | `dst, src0, src1` | `docs/isa/TREM.md` |67| Elementwise (Tile-Tile) | `TREM` | `elementwise` | `dst, src0, src1` | `docs/isa/TREM.md` |
68| Elementwise (Tile-Tile) | `TFMOD` | `elementwise` | `dst, src0, src1` | `docs/isa/TFMOD.md` |68| Elementwise (Tile-Tile) | `TFMOD` | `elementwise` | `dst, src0, src1` | `docs/isa/TFMOD.md` |
69-| Elementwise (Tile-Tile) | `TMULADDDST` | `elementwise` | `dst, src0, src1` | `docs/isa/TMULADDDST.md` |69+| Elementwise (Tile-Tile) | `TMULA` | `elementwise` | `dst, src0, src1` | `docs/isa/TMULA.md` |
70| Tile-Scalar / Tile-Immediate | `TEXPANDS` | `scalar` | `dst, src, scalar` | `docs/isa/TEXPANDS.md` |70| Tile-Scalar / Tile-Immediate | `TEXPANDS` | `scalar` | `dst, src, scalar` | `docs/isa/TEXPANDS.md` |
71| Tile-Scalar / Tile-Immediate | `TCMPS` | `scalar` | `dst, src, scalar` | `docs/isa/TCMPS.md` |71| Tile-Scalar / Tile-Immediate | `TCMPS` | `scalar` | `dst, src, scalar` | `docs/isa/TCMPS.md` |
72| Tile-Scalar / Tile-Immediate | `TSELS` | `scalar` | `dst, src, scalar` | `docs/isa/TSELS.md` |72| Tile-Scalar / Tile-Immediate | `TSELS` | `scalar` | `dst, src, scalar` | `docs/isa/TSELS.md` |
Mdocs/mkdocs/src/manual/appendix-d-instruction-family-matrix_zh.md+1-1
@@ -66,7 +66,7 @@
66| 逐元素（Tile-Tile） | `TNEG` | `elementwise` | `dst, src0, src1` | `docs/isa/TNEG_zh.md` |66| 逐元素（Tile-Tile） | `TNEG` | `elementwise` | `dst, src0, src1` | `docs/isa/TNEG_zh.md` |
67| 逐元素（Tile-Tile） | `TREM` | `elementwise` | `dst, src0, src1` | `docs/isa/TREM_zh.md` |67| 逐元素（Tile-Tile） | `TREM` | `elementwise` | `dst, src0, src1` | `docs/isa/TREM_zh.md` |
68| 逐元素（Tile-Tile） | `TFMOD` | `elementwise` | `dst, src0, src1` | `docs/isa/TFMOD_zh.md` |68| 逐元素（Tile-Tile） | `TFMOD` | `elementwise` | `dst, src0, src1` | `docs/isa/TFMOD_zh.md` |
69-| 逐元素（Tile-Tile） | `TMULADDDST` | `elementwise` | `dst, src0, src1` | `docs/isa/TMULADDDST_zh.md` |69+| 逐元素（Tile-Tile） | `TMULA` | `elementwise` | `dst, src0, src1` | `docs/isa/TMULA_zh.md` |
70| Tile-标量 / Tile-立即数 | `TEXPANDS` | `scalar` | `dst, src, scalar` | `docs/isa/TEXPANDS_zh.md` |70| Tile-标量 / Tile-立即数 | `TEXPANDS` | `scalar` | `dst, src, scalar` | `docs/isa/TEXPANDS_zh.md` |
71| Tile-标量 / Tile-立即数 | `TCMPS` | `scalar` | `dst, src, scalar` | `docs/isa/TCMPS_zh.md` |71| Tile-标量 / Tile-立即数 | `TCMPS` | `scalar` | `dst, src, scalar` | `docs/isa/TCMPS_zh.md` |
72| Tile-标量 / Tile-立即数 | `TSELS` | `scalar` | `dst, src, scalar` | `docs/isa/TSELS_zh.md` |72| Tile-标量 / Tile-立即数 | `TSELS` | `scalar` | `dst, src, scalar` | `docs/isa/TSELS_zh.md` |
Mdocs/tools/gen_isa_indexes.py+2-2
@@ -205,7 +205,7 @@ def append_removed_interfaces_en(lines: List[str]) -> None:
205 "- Replace ternary/scalar fused arithmetic forms with the corresponding primitive arithmetic sequence, "205 "- Replace ternary/scalar fused arithmetic forms with the corresponding primitive arithmetic sequence, "
206 "such as `TADD`, `TSUB`, `TADDS`, `TSUBS`, `TMUL`, `TMADD`, and `TRELU`."206 "such as `TADD`, `TSUB`, `TADDS`, `TSUBS`, `TMUL`, `TMADD`, and `TRELU`."
207 )207 )
208- lines.append("- Rename legacy fused multiply-add APIs: `TFUSEDMULADD` to `TMADD`.")208+ lines.append("- Rename legacy fused multiply-add APIs: `TFUSEDMULADD` to `TMADD`, and `TMULADDDST` to `TMULA`.")
209 lines.append(209 lines.append(
210 "- Replace fused add/ReLU/convert or add/dequant/ReLU forms with explicit arithmetic, "210 "- Replace fused add/ReLU/convert or add/dequant/ReLU forms with explicit arithmetic, "
211 "conversion/dequantization, and `TRELU` steps."211 "conversion/dequantization, and `TRELU` steps."
@@ -245,7 +245,7 @@ def append_removed_interfaces_zh(lines: List[str]) -> None:
245 "- 将三元/标量融合算术形式替换为对应基础算术序列，例如 `TADD`、`TSUB`、`TADDS`、`TSUBS`、"245 "- 将三元/标量融合算术形式替换为对应基础算术序列，例如 `TADD`、`TSUB`、`TADDS`、`TSUBS`、"
246 "`TMUL`、`TMADD` 和 `TRELU`。"246 "`TMUL`、`TMADD` 和 `TRELU`。"
247 )247 )
248- lines.append("- 将历史融合乘加接口重命名：`TFUSEDMULADD` 改为 `TMADD`。")248+ lines.append("- 将历史融合乘加接口重命名：`TFUSEDMULADD` 改为 `TMADD`，`TMULADDDST` 改为 `TMULA`。")
249 lines.append("- 将融合 add/ReLU/convert 或 add/dequant/ReLU 形式拆分为显式算术、转换/反量化和 `TRELU` 步骤。")249 lines.append("- 将融合 add/ReLU/convert 或 add/dequant/ReLU 形式拆分为显式算术、转换/反量化和 `TRELU` 步骤。")
250 lines.append("- 将 `TPairReduceSum` 替换为与目标 layout 匹配的现有行/列归约原语。")250 lines.append("- 将 `TPairReduceSum` 替换为与目标 layout 匹配的现有行/列归约原语。")
251 lines.append(251 lines.append(
Mdocs/tools/gen_isa_svg_core.py+1-1
@@ -629,7 +629,7 @@ BINARY_ELEMENTWISE_EXPR = {
629 629 
630TERNARY_ELEMENTWISE_SPEC = {630TERNARY_ELEMENTWISE_SPEC = {
631 "TMADD": ("src0 * dst(old) + src1", ["src0", "dst(old)", "src1"]),631 "TMADD": ("src0 * dst(old) + src1", ["src0", "dst(old)", "src1"]),
632- "TMULADDDST": ("src0 * src1 + dst(old)", ["src0", "src1", "dst(old)"]),632+ "TMULA": ("src0 * src1 + dst(old)", ["src0", "src1", "dst(old)"]),
633}633}
634 634 
635 635 
Minclude/README.md+1-1
@@ -167,7 +167,7 @@ This table tracks per-instruction backend availability:
167| [`TWAIT`](../docs/isa/comm/TWAIT.md) | Yes | TODO | Yes | Yes | Yes | TODO |167| [`TWAIT`](../docs/isa/comm/TWAIT.md) | Yes | TODO | Yes | Yes | Yes | TODO |
168| [`TXOR`](../docs/isa/TXOR.md) | Yes | TODO | Yes | Yes | Yes | Yes |168| [`TXOR`](../docs/isa/TXOR.md) | Yes | TODO | Yes | Yes | Yes | Yes |
169| [`TXORS`](../docs/isa/TXORS.md) | Yes | TODO | Yes | Yes | Yes | Yes |169| [`TXORS`](../docs/isa/TXORS.md) | Yes | TODO | Yes | Yes | Yes | Yes |
170-| [`TMULADDDST`](../docs/isa/TMULADDDST_zh.md) | TODO | TODO | Yes | Yes | Yes | TODO |170+| [`TMULA`](../docs/isa/TMULA_zh.md) | TODO | TODO | Yes | Yes | Yes | TODO |
171 171 
172Notes:172Notes:
173 173 
Minclude/README_zh.md+1-1
@@ -166,7 +166,7 @@ PTO Tile Lib 对外的 C/C++ 头文件（以模板化、基本 header-only 为
166| [`TWAIT`](../docs/isa/comm/TWAIT_zh.md) | 是 | TODO | 是 | 是 | 是 | TODO |166| [`TWAIT`](../docs/isa/comm/TWAIT_zh.md) | 是 | TODO | 是 | 是 | 是 | TODO |
167| [`TXOR`](../docs/isa/TXOR_zh.md) | 是 | TODO | 是 | 是 | 是 | 是 |167| [`TXOR`](../docs/isa/TXOR_zh.md) | 是 | TODO | 是 | 是 | 是 | 是 |
168| [`TXORS`](../docs/isa/TXORS_zh.md) | 是 | TODO | 是 | 是 | 是 | 是 |168| [`TXORS`](../docs/isa/TXORS_zh.md) | 是 | TODO | 是 | 是 | 是 | 是 |
169-| [`TMULADDDST`](../docs/isa/TMULADDDST_zh.md) | TODO | TODO | 是 | 是 | 是 | TODO |169+| [`TMULA`](../docs/isa/TMULA_zh.md) | TODO | TODO | 是 | 是 | 是 | TODO |
170 170 
171说明：171说明：
172 172 
Minclude/pto/common/event.hpp+2-2
@@ -137,7 +137,7 @@ enum class Op : uint16_t {
137 TCOLARGMIN,137 TCOLARGMIN,
138 TROWARGMAX,138 TROWARGMAX,
139 TROWARGMIN,139 TROWARGMIN,
140- TMULADDDST,140+ TMULA,
141 TPREFETCH,141 TPREFETCH,
142 TFILLPAD_VEC,142 TFILLPAD_VEC,
143 TFILLPAD_MAT,143 TFILLPAD_MAT,
@@ -286,7 +286,7 @@ PTO_DEFINE_OP_PIPE(Op::TFILLPAD_MAT, PIPE_MTE2);
286PTO_DEFINE_OP_PIPE(Op::MGATHER_VEC, PIPE_V);286PTO_DEFINE_OP_PIPE(Op::MGATHER_VEC, PIPE_V);
287PTO_DEFINE_OP_PIPE(Op::MGATHER_MAT, PIPE_MTE2);287PTO_DEFINE_OP_PIPE(Op::MGATHER_MAT, PIPE_MTE2);
288PTO_DEFINE_OP_PIPE(Op::MSCATTER, PIPE_V);288PTO_DEFINE_OP_PIPE(Op::MSCATTER, PIPE_V);
289-PTO_DEFINE_OP_PIPE(Op::TMULADDDST, PIPE_V);289+PTO_DEFINE_OP_PIPE(Op::TMULA, PIPE_V);
290PTO_DEFINE_OP_PIPE(Op::TQUANT, PIPE_V);290PTO_DEFINE_OP_PIPE(Op::TQUANT, PIPE_V);
291PTO_DEFINE_OP_PIPE(Op::THISTOGRAM, PIPE_V);291PTO_DEFINE_OP_PIPE(Op::THISTOGRAM, PIPE_V);
292PTO_DEFINE_OP_PIPE(Op::TINTERLEAVE, PIPE_V);292PTO_DEFINE_OP_PIPE(Op::TINTERLEAVE, PIPE_V);
Minclude/pto/common/pto_instr.hpp+2-2
@@ -1400,10 +1400,10 @@ PTO_INST RecordEvent TPARTARGMIN(
1400}1400}
1401 1401 
1402template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>1402template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>
1403-PTO_INST RecordEvent TMULADDDST(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1, WaitEvents&... events)1403+PTO_INST RecordEvent TMULA(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1, WaitEvents&... events)
1404{1404{
1405 detail::PtoWaitEvents(events...);1405 detail::PtoWaitEvents(events...);
1406- MAP_INSTR_IMPL(TMULADDDST, dst, src0, src1);1406+ MAP_INSTR_IMPL(TMULA, dst, src0, src1);
1407 return {};1407 return {};
1408}1408}
1409 1409 
Minclude/pto/common/pto_instr_impl.hpp+2-2
@@ -184,7 +184,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
184#include "pto/npu/a2a3/TAlloc.hpp"184#include "pto/npu/a2a3/TAlloc.hpp"
185#include "pto/npu/a2a3/TFree.hpp"185#include "pto/npu/a2a3/TFree.hpp"
186#include "pto/npu/a2a3/TColReduceIdx.hpp"186#include "pto/npu/a2a3/TColReduceIdx.hpp"
187-#include "pto/npu/a2a3/TMulAddDst.hpp"187+#include "pto/npu/a2a3/TMula.hpp"
188#endif188#endif
189#endif189#endif
190 190 
@@ -310,7 +310,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
310#include "pto/npu/a5/TColReduceIdx.hpp"310#include "pto/npu/a5/TColReduceIdx.hpp"
311#include "pto/npu/a5/TInterleave.hpp"311#include "pto/npu/a5/TInterleave.hpp"
312#include "pto/npu/a5/TDeInterleave.hpp"312#include "pto/npu/a5/TDeInterleave.hpp"
313-#include "pto/npu/a5/TMulAddDst.hpp"313+#include "pto/npu/a5/TMula.hpp"
314#endif // __COSTMODEL314#endif // __COSTMODEL
315#endif315#endif
316 316 
Minclude/pto/cpu/ElementOp.h+2-2
@@ -26,7 +26,7 @@ enum class ElementOp {
26 OP_SUB,26 OP_SUB,
27 OP_MUL,27 OP_MUL,
28 OP_DIV,28 OP_DIV,
29- OP_MULADDDST,29+ OP_MULA,
30 OP_MADD,30 OP_MADD,
31 OP_REM,31 OP_REM,
32 OP_SHL,32 OP_SHL,
@@ -133,7 +133,7 @@ struct ElementOpCal<DType, ElementOp::OP_DIV> {
133};133};
134 134 
135template <typename DType>135template <typename DType>
136-struct ElementOpCal<DType, ElementOp::OP_MULADDDST> {136+struct ElementOpCal<DType, ElementOp::OP_MULA> {
137 static void apply(DType& dst, DType& src0, DType& src1, size_t) { dst = static_cast<DType>(src0 * src1) + dst; }137 static void apply(DType& dst, DType& src0, DType& src1, size_t) { dst = static_cast<DType>(src0 * src1) + dst; }
138 138 
139 static void apply(DType& dst, const DType& src0, const DType& src1)139 static void apply(DType& dst, const DType& src0, const DType& src1)
Minclude/pto/cpu/ElementTileOp.h+1-1
@@ -114,7 +114,7 @@ BINARY_OP_DEF(AND)
114BINARY_OP_DEF(OR)114BINARY_OP_DEF(OR)
115BINARY_OP_DEF(XOR)115BINARY_OP_DEF(XOR)
116BINARY_OP_DEF(MIN)116BINARY_OP_DEF(MIN)
117-BINARY_OP_DEF(MULADDDST)117+BINARY_OP_DEF(MULA)
118BINARY_OP_DEF(MADD)118BINARY_OP_DEF(MADD)
119 119 
120UNARY_OP_DEF(NEG)120UNARY_OP_DEF(NEG)
Rinclude/pto/npu/a2a3/TMulAddDst.hpp→include/pto/npu/a2a3/TMula.hpp+14-14
@@ -8,8 +8,8 @@ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A
8See LICENSE in the root of the software repository for the full text of the License.8See LICENSE in the root of the software repository for the full text of the License.
9*/9*/
10 10 
11-#ifndef TMULADDDST_HPP11+#ifndef TMULA_HPP
12-#define TMULADDDST_HPP12+#define TMULA_HPP
13 13 
14#include <pto/common/constants.hpp>14#include <pto/common/constants.hpp>
15#include <pto/common/utils.hpp>15#include <pto/common/utils.hpp>
@@ -18,7 +18,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
18namespace pto {18namespace pto {
19 19 
20template <typename T>20template <typename T>
21-struct MulAddDstOp {21+struct MulaOp {
22 PTO_INTERNAL static void BinInstr(__ubuf__ T* dst, __ubuf__ T* src0, __ubuf__ T* src1, uint8_t repeats)22 PTO_INTERNAL static void BinInstr(__ubuf__ T* dst, __ubuf__ T* src0, __ubuf__ T* src1, uint8_t repeats)
23 {23 {
24 vmla(dst, src0, src1, repeats, 1, 1, 1, 8, 8, 8);24 vmla(dst, src0, src1, repeats, 1, 1, 1, 8, 8, 8);
@@ -34,7 +34,7 @@ struct MulAddDstOp {
34template <34template <
35 typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, unsigned elementsPerRepeat,35 typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, unsigned elementsPerRepeat,
36 unsigned blockSizeElem, unsigned dstRowStride, unsigned src0RowStride, unsigned src1RowStride>36 unsigned blockSizeElem, unsigned dstRowStride, unsigned src0RowStride, unsigned src1RowStride>
37-__tf__ PTO_INTERNAL void TMulAddDst(37+__tf__ PTO_INTERNAL void TMula(
38 typename TileDataDst::TileDType __in__ __out__ dst, typename TileDataSrc0::TileDType __in__ src0,38 typename TileDataDst::TileDType __in__ __out__ dst, typename TileDataSrc0::TileDType __in__ src0,
39 typename TileDataSrc1::TileDType __in__ src1, unsigned validRows, unsigned validCols)39 typename TileDataSrc1::TileDType __in__ src1, unsigned validRows, unsigned validCols)
40{40{
@@ -43,49 +43,49 @@ __tf__ PTO_INTERNAL void TMulAddDst(
43 __ubuf__ T* src0Ptr = (__ubuf__ T*)__cce_get_tile_ptr(src0);43 __ubuf__ T* src0Ptr = (__ubuf__ T*)__cce_get_tile_ptr(src0);
44 __ubuf__ T* src1Ptr = (__ubuf__ T*)__cce_get_tile_ptr(src1);44 __ubuf__ T* src1Ptr = (__ubuf__ T*)__cce_get_tile_ptr(src1);
45 if constexpr (dstRowStride == src0RowStride && dstRowStride == src1RowStride) {45 if constexpr (dstRowStride == src0RowStride && dstRowStride == src1RowStride) {
46- BinaryInstr<MulAddDstOp<T>, T, TileDataDst, elementsPerRepeat, blockSizeElem, dstRowStride>(46+ BinaryInstr<MulaOp<T>, T, TileDataDst, elementsPerRepeat, blockSizeElem, dstRowStride>(
47 dstPtr, src0Ptr, src1Ptr, validRows, validCols);47 dstPtr, src0Ptr, src1Ptr, validRows, validCols);
48 } else {48 } else {
49- BinaryInstr<MulAddDstOp<T>, T, elementsPerRepeat, blockSizeElem, dstRowStride, src0RowStride, src1RowStride>(49+ BinaryInstr<MulaOp<T>, T, elementsPerRepeat, blockSizeElem, dstRowStride, src0RowStride, src1RowStride>(
50 dstPtr, src0Ptr, src1Ptr, validRows, validCols);50 dstPtr, src0Ptr, src1Ptr, validRows, validCols);
51 }51 }
52 return;52 return;
53}53}
54 54 
55template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>55template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>
56-PTO_INTERNAL void TMulAddDstCheck(const TileDataDst& dst, const TileDataSrc0& src0, const TileDataSrc1& src1)56+PTO_INTERNAL void TMulaCheck(const TileDataDst& dst, const TileDataSrc0& src0, const TileDataSrc1& src1)
57{57{
58 using T = typename TileDataDst::DType;58 using T = typename TileDataDst::DType;
59 static_assert(59 static_assert(
60 std::is_same_v<T, typename TileDataSrc0::DType> && std::is_same_v<T, typename TileDataSrc1::DType>,60 std::is_same_v<T, typename TileDataSrc0::DType> && std::is_same_v<T, typename TileDataSrc1::DType>,
61- "Fix: TMULADDDST the data type of dst must be consistent with of src0 and src1.");61+ "Fix: TMULA the data type of dst must be consistent with of src0 and src1.");
62 static_assert(62 static_assert(
63 std::is_same_v<T, half> || std::is_same_v<T, float16_t> || std::is_same_v<T, float> ||63 std::is_same_v<T, half> || std::is_same_v<T, float16_t> || std::is_same_v<T, float> ||
64 std::is_same_v<T, float32_t>,64 std::is_same_v<T, float32_t>,
65- "Fix: TMULADDDST has invalid data type.");65+ "Fix: TMULA has invalid data type.");
66 static_assert(66 static_assert(
67 TileDataDst::isRowMajor && TileDataSrc0::isRowMajor && TileDataSrc1::isRowMajor,67 TileDataDst::isRowMajor && TileDataSrc0::isRowMajor && TileDataSrc1::isRowMajor,
68- "Fix: TMULADDDST only support row major layout.");68+ "Fix: TMULA only support row major layout.");
69 unsigned validRows = dst.GetValidRow();69 unsigned validRows = dst.GetValidRow();
70 unsigned validCols = dst.GetValidCol();70 unsigned validCols = dst.GetValidCol();
71 PTO_ASSERT(71 PTO_ASSERT(
72 src0.GetValidRow() == validRows && src0.GetValidCol() == validCols && src1.GetValidRow() == validRows &&72 src0.GetValidRow() == validRows && src0.GetValidCol() == validCols && src1.GetValidRow() == validRows &&
73 src1.GetValidCol() == validCols,73 src1.GetValidCol() == validCols,
74- "Fix: TMULADDDST input tile src0 valid shape mismatch with output tile dst shape.");74+ "Fix: TMULA input tile src0 valid shape mismatch with output tile dst shape.");
75}75}
76 76 
77template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>77template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>
78-PTO_INTERNAL void TMULADDDST_IMPL(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1)78+PTO_INTERNAL void TMULA_IMPL(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1)
79{79{
80 using T = typename TileDataDst::DType;80 using T = typename TileDataDst::DType;
81- TMulAddDstCheck<TileDataDst, TileDataSrc0, TileDataSrc1>(dst, src0, src1);81+ TMulaCheck<TileDataDst, TileDataSrc0, TileDataSrc1>(dst, src0, src1);
82 constexpr unsigned blockSizeElem = BLOCK_BYTE_SIZE / sizeof(T);82 constexpr unsigned blockSizeElem = BLOCK_BYTE_SIZE / sizeof(T);
83 constexpr unsigned elementsPerRepeat = REPEAT_BYTE / sizeof(T);83 constexpr unsigned elementsPerRepeat = REPEAT_BYTE / sizeof(T);
84 84 
85 constexpr unsigned dstRowStride = TileDataDst::RowStride;85 constexpr unsigned dstRowStride = TileDataDst::RowStride;
86 constexpr unsigned src0RowStride = TileDataSrc0::RowStride;86 constexpr unsigned src0RowStride = TileDataSrc0::RowStride;
87 constexpr unsigned src1RowStride = TileDataSrc1::RowStride;87 constexpr unsigned src1RowStride = TileDataSrc1::RowStride;
88- TMulAddDst<88+ TMula<
89 TileDataDst, TileDataSrc0, TileDataSrc1, elementsPerRepeat, blockSizeElem, dstRowStride, src0RowStride,89 TileDataDst, TileDataSrc0, TileDataSrc1, elementsPerRepeat, blockSizeElem, dstRowStride, src0RowStride,
90 src1RowStride>(dst.data(), src0.data(), src1.data(), dst.GetValidRow(), dst.GetValidCol());90 src1RowStride>(dst.data(), src0.data(), src1.data(), dst.GetValidRow(), dst.GetValidCol());
91}91}
Rinclude/pto/npu/a5/TMulAddDst.hpp→include/pto/npu/a5/TMula.hpp+13-13
@@ -8,8 +8,8 @@ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A
8See LICENSE in the root of the software repository for the full text of the License.8See LICENSE in the root of the software repository for the full text of the License.
9*/9*/
10 10 
11-#ifndef TMULADDDST_HPP11+#ifndef TMULA_HPP
12-#define TMULADDDST_HPP12+#define TMULA_HPP
13 13 
14#include <pto/common/constants.hpp>14#include <pto/common/constants.hpp>
15#include <pto/common/utils.hpp>15#include <pto/common/utils.hpp>
@@ -20,7 +20,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
20namespace pto {20namespace pto {
21 21 
22template <typename T>22template <typename T>
23-struct MulAddDstOp {23+struct MulaOp {
24 PTO_INTERNAL static void TernInstr(24 PTO_INTERNAL static void TernInstr(
25 RegTensor<T>& reg_dst, RegTensor<T>& reg_src0, RegTensor<T>& reg_src1, MaskReg& preg)25 RegTensor<T>& reg_dst, RegTensor<T>& reg_src0, RegTensor<T>& reg_src1, MaskReg& preg)
26 {26 {
@@ -36,7 +36,7 @@ struct MulAddDstOp {
36template <36template <
37 typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, unsigned ElementsPerRepeat,37 typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, unsigned ElementsPerRepeat,
38 unsigned BlockSizeElem>38 unsigned BlockSizeElem>
39-__tf__ PTO_INTERNAL OP_NAME(TMULADDDST) OP_TYPE(element_wise) void TMulAddDst(39+__tf__ PTO_INTERNAL OP_NAME(TMULA) OP_TYPE(element_wise) void TMula(
40 typename TileDataDst::TileDType __out__ dst, typename TileDataSrc0::TileDType __in__ src0,40 typename TileDataDst::TileDType __out__ dst, typename TileDataSrc0::TileDType __in__ src0,
41 typename TileDataSrc1::TileDType __in__ src1, unsigned validRows, unsigned validCols,41 typename TileDataSrc1::TileDType __in__ src1, unsigned validRows, unsigned validCols,
42 VFImplKind version = VFImplKind::VFIMPL_DEFAULT)42 VFImplKind version = VFImplKind::VFIMPL_DEFAULT)
@@ -46,42 +46,42 @@ __tf__ PTO_INTERNAL OP_NAME(TMULADDDST) OP_TYPE(element_wise) void TMulAddDst(
46 __ubuf__ T* src0Ptr = (__ubuf__ T*)__cce_get_tile_ptr(src0);46 __ubuf__ T* src0Ptr = (__ubuf__ T*)__cce_get_tile_ptr(src0);
47 __ubuf__ T* src1Ptr = (__ubuf__ T*)__cce_get_tile_ptr(src1);47 __ubuf__ T* src1Ptr = (__ubuf__ T*)__cce_get_tile_ptr(src1);
48 48 
49- TernaryInstr<MulAddDstOp<T>, TileDataDst, TileDataSrc0, TileDataSrc1, ElementsPerRepeat, BlockSizeElem>(49+ TernaryInstr<MulaOp<T>, TileDataDst, TileDataSrc0, TileDataSrc1, ElementsPerRepeat, BlockSizeElem>(
50 dstPtr, src0Ptr, src1Ptr, validRows, validCols, version);50 dstPtr, src0Ptr, src1Ptr, validRows, validCols, version);
51 return;51 return;
52}52}
53 53 
54template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>54template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>
55-PTO_INTERNAL void TMulAddDstCheck(const TileDataDst& dst, const TileDataSrc0& src0, const TileDataSrc1& src1)55+PTO_INTERNAL void TMulaCheck(const TileDataDst& dst, const TileDataSrc0& src0, const TileDataSrc1& src1)
56{56{
57 unsigned validRows = dst.GetValidRow();57 unsigned validRows = dst.GetValidRow();
58 unsigned validCols = dst.GetValidCol();58 unsigned validCols = dst.GetValidCol();
59 PTO_ASSERT(59 PTO_ASSERT(
60 src0.GetValidRow() == validRows && src0.GetValidCol() == validCols && src1.GetValidRow() == validRows &&60 src0.GetValidRow() == validRows && src0.GetValidCol() == validCols && src1.GetValidRow() == validRows &&
61 src1.GetValidCol() == validCols,61 src1.GetValidCol() == validCols,
62- "Fix: TMULADDDST input tile src0 valid shape mismatch with output tile dst shape.");62+ "Fix: TMULA input tile src0 valid shape mismatch with output tile dst shape.");
63 using T = typename TileDataDst::DType;63 using T = typename TileDataDst::DType;
64 static_assert(64 static_assert(
65 TileDataDst::isRowMajor && TileDataSrc0::isRowMajor && TileDataSrc1::isRowMajor,65 TileDataDst::isRowMajor && TileDataSrc0::isRowMajor && TileDataSrc1::isRowMajor,
66- "Fix: TMULADDDST only support row major layout.");66+ "Fix: TMULA only support row major layout.");
67 static_assert(67 static_assert(
68 std::is_same_v<T, typename TileDataSrc0::DType> && std::is_same_v<T, typename TileDataSrc1::DType>,68 std::is_same_v<T, typename TileDataSrc0::DType> && std::is_same_v<T, typename TileDataSrc1::DType>,
69- "Fix: TMULADDDST the data type of dst must be consistent with of src0 and src1.");69+ "Fix: TMULA the data type of dst must be consistent with of src0 and src1.");
70 static_assert(70 static_assert(
71 std::is_same_v<T, half> || std::is_same_v<T, float16_t> || std::is_same_v<T, float> ||71 std::is_same_v<T, half> || std::is_same_v<T, float16_t> || std::is_same_v<T, float> ||
72 std::is_same_v<T, float32_t>,72 std::is_same_v<T, float32_t>,
73- "Fix: TMULADDDST has invalid data type.");73+ "Fix: TMULA has invalid data type.");
74}74}
75 75 
76template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>76template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>
77-PTO_INTERNAL void TMULADDDST_IMPL(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1)77+PTO_INTERNAL void TMULA_IMPL(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1)
78{78{
79 using T = typename TileDataDst::DType;79 using T = typename TileDataDst::DType;
80- TMulAddDstCheck<TileDataDst, TileDataSrc0, TileDataSrc1>(dst, src0, src1);80+ TMulaCheck<TileDataDst, TileDataSrc0, TileDataSrc1>(dst, src0, src1);
81 constexpr unsigned blockSizeElem = BLOCK_BYTE_SIZE / sizeof(T);81 constexpr unsigned blockSizeElem = BLOCK_BYTE_SIZE / sizeof(T);
82 constexpr unsigned elementsPerRepeat = CCE_VL / sizeof(T);82 constexpr unsigned elementsPerRepeat = CCE_VL / sizeof(T);
83 83 
84- TMulAddDst<TileDataDst, TileDataSrc0, TileDataSrc1, elementsPerRepeat, blockSizeElem>(84+ TMula<TileDataDst, TileDataSrc0, TileDataSrc1, elementsPerRepeat, blockSizeElem>(
85 dst.data(), src0.data(), src1.data(), dst.GetValidRow(), dst.GetValidCol());85 dst.data(), src0.data(), src1.data(), dst.GetValidRow(), dst.GetValidCol());
86}86}
87} // namespace pto87} // namespace pto
Mtests/cpu/st/testcase/CMakeLists.txt+1-1
@@ -100,7 +100,7 @@ tmins
100tmov100tmov
101tmrgsort101tmrgsort
102tmul102tmul
103-tmuladddst103+tmula
104tmuls104tmuls
105tneg105tneg
106tnot106tnot
Rtests/npu/a2a3/src/st/testcase/tmuladddst/CMakeLists.txt→tests/cpu/st/testcase/tmula/CMakeLists.txt+1-1
@@ -8,4 +8,4 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# --------------------------------------------------------------------------------9# --------------------------------------------------------------------------------
10 10 
11-pto_vec_st(tmuladddst)11+pto_cpu_sim_st(tmula)
Rtests/cpu/st/testcase/tmuladddst/gen_data.py→tests/cpu/st/testcase/tmula/gen_data.py+1-1
@@ -90,7 +90,7 @@ class TestParams:
90 np.uint16: 'uint16',90 np.uint16: 'uint16',
91 np.uint8: 'uint8'91 np.uint8: 'uint8'
92 }[dtype]92 }[dtype]
93- self.name = f"TMULADDDSTTest.case_{dtype_str}_{dst_tile_row}x{dst_tile_col}_\93+ self.name = f"TMULATest.case_{dtype_str}_{dst_tile_row}x{dst_tile_col}_\
94{src0_tile_row}x{src0_tile_col}_{src1_tile_row}x{src1_tile_col}_\94{src0_tile_row}x{src0_tile_col}_{src1_tile_row}x{src1_tile_col}_\
95{valid_row}x{valid_col}"95{valid_row}x{valid_col}"
96 96 
Rtests/cpu/st/testcase/tmuladddst/main.cpp→tests/cpu/st/testcase/tmula/main.cpp+11-17
@@ -15,7 +15,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
15using namespace std;15using namespace std;
16using namespace PtoTestCommon;16using namespace PtoTestCommon;
17 17 
18-class TMULADDDSTTest : public testing::Test {18+class TMULATest : public testing::Test {
19protected:19protected:
20 void SetUp() override {}20 void SetUp() override {}
21 void TearDown() override {}21 void TearDown() override {}
@@ -33,12 +33,12 @@ std::string GetGoldenDir()
33template <33template <
34 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,34 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
35 int vCols, bool isHalf = true>35 int vCols, bool isHalf = true>
36-void LaunchTMULADDDST(T* out, T* src0, T* src1, void* stream);36+void LaunchTMULA(T* out, T* src0, T* src1, void* stream);
37 37 
38template <38template <
39 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,39 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
40 int vCols, bool isHalf = false>40 int vCols, bool isHalf = false>
41-void test_TMULADDDST()41+void test_TMULA()
42{42{
43 size_t fileSizeDst = dstTileH * dstTileW * sizeof(T);43 size_t fileSizeDst = dstTileH * dstTileW * sizeof(T);
44 size_t fileSizeSrc0 = src0TileH * src0TileW * sizeof(T);44 size_t fileSizeSrc0 = src0TileH * src0TileW * sizeof(T);
@@ -68,7 +68,7 @@ void test_TMULADDDST()
68 aclrtMemcpy(src0Device, fileSizeSrc0, src0Host, fileSizeSrc0, ACL_MEMCPY_HOST_TO_DEVICE);68 aclrtMemcpy(src0Device, fileSizeSrc0, src0Host, fileSizeSrc0, ACL_MEMCPY_HOST_TO_DEVICE);
69 aclrtMemcpy(src1Device, fileSizeSrc1, src1Host, fileSizeSrc1, ACL_MEMCPY_HOST_TO_DEVICE);69 aclrtMemcpy(src1Device, fileSizeSrc1, src1Host, fileSizeSrc1, ACL_MEMCPY_HOST_TO_DEVICE);
70 70 
71- LaunchTMULADDDST<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(71+ LaunchTMULA<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(
72 dstDevice, src0Device, src1Device, stream);72 dstDevice, src0Device, src1Device, stream);
73 73 
74 aclrtSynchronizeStream(stream);74 aclrtSynchronizeStream(stream);
@@ -97,21 +97,15 @@ void test_TMULADDDST()
97 EXPECT_TRUE(ret);97 EXPECT_TRUE(ret);
98}98}
99 99 
100-TEST_F(TMULADDDSTTest, case_float_64x64_64x64_64x64_64x64) { test_TMULADDDST<float, 64, 64, 64, 64, 64, 64, 64, 64>(); }100+TEST_F(TMULATest, case_float_64x64_64x64_64x64_64x64) { test_TMULA<float, 64, 64, 64, 64, 64, 64, 64, 64>(); }
101-TEST_F(TMULADDDSTTest, case_float_32x128_32x192_32x256_32x127)101+TEST_F(TMULATest, case_float_32x128_32x192_32x256_32x127) { test_TMULA<float, 32, 128, 32, 192, 32, 256, 32, 127>(); }
102+TEST_F(TMULATest, case_half_64x64_64x64_64x64_64x64) { test_TMULA<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(); }
103+TEST_F(TMULATest, case_half_32x128_32x192_32x256_32x127)
102{104{
103- test_TMULADDDST<float, 32, 128, 32, 192, 32, 256, 32, 127>();105+ test_TMULA<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>();
104-}
105-TEST_F(TMULADDDSTTest, case_half_64x64_64x64_64x64_64x64)
106-{
107- test_TMULADDDST<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>();
108-}
109-TEST_F(TMULADDDSTTest, case_half_32x128_32x192_32x256_32x127)
110-{
111- test_TMULADDDST<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>();
112}106}
113 107 
114-TEST_F(TMULADDDSTTest, case_half_1x16384_1x16384_1x16384_1x16384)108+TEST_F(TMULATest, case_half_1x16384_1x16384_1x16384_1x16384)
115{109{
116- test_TMULADDDST<aclFloat16, 1, 16384, 1, 16384, 1, 16384, 1, 16384>();110+ test_TMULA<aclFloat16, 1, 16384, 1, 16384, 1, 16384, 1, 16384>();
117}111}
Rtests/cpu/st/testcase/tmuladddst/tmuladddst_kernel.cpp→tests/cpu/st/testcase/tmula/tmula_kernel.cpp+10-11
@@ -15,7 +15,7 @@ using namespace pto;
15template <15template <
16 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,16 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
17 int vCols>17 int vCols>
18-__global__ AICORE void runTMULADDDST(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1)18+__global__ AICORE void runTMULA(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1)
19{19{
20 using DynShape = pto::Shape<-1, -1, -1, -1, -1>;20 using DynShape = pto::Shape<-1, -1, -1, -1, -1>;
21 using DynStride = pto::Stride<-1, -1, -1, -1, -1>;21 using DynStride = pto::Stride<-1, -1, -1, -1, -1>;
@@ -47,7 +47,7 @@ __global__ AICORE void runTMULADDDST(__gm__ T __out__* out, __gm__ T __in__* src
47 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);47 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
48 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);48 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
49#endif49#endif
50- TMULADDDST<TileDataDst, TileDataSrc0, TileDataSrc1>(dstTile, src0Tile, src1Tile);50+ TMULA<TileDataDst, TileDataSrc0, TileDataSrc1>(dstTile, src0Tile, src1Tile);
51#ifndef __PTO_AUTO__51#ifndef __PTO_AUTO__
52 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);52 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
53 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);53 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
@@ -59,23 +59,22 @@ __global__ AICORE void runTMULADDDST(__gm__ T __out__* out, __gm__ T __in__* src
59template <59template <
60 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,60 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
61 int vCols, bool isHalf = true>61 int vCols, bool isHalf = true>
62-void LaunchTMULADDDST(T* out, T* src0, T* src1, void* stream)62+void LaunchTMULA(T* out, T* src0, T* src1, void* stream)
63{63{
64 if constexpr (std::is_same_v<T, aclFloat16>) {64 if constexpr (std::is_same_v<T, aclFloat16>) {
65- runTMULADDDST<half, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(65+ runTMULA<half, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(
66 (half*)(out), (half*)(src0), (half*)(src1));66 (half*)(out), (half*)(src0), (half*)(src1));
67 } else {67 } else {
68- runTMULADDDST<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(out, src0, src1);68+ runTMULA<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(out, src0, src1);
69 }69 }
70}70}
71 71 
72-template void LaunchTMULADDDST<float, 64, 64, 64, 64, 64, 64, 64, 64>(72+template void LaunchTMULA<float, 64, 64, 64, 64, 64, 64, 64, 64>(float* out, float* src0, float* src1, void* stream);
73+template void LaunchTMULA<float, 32, 128, 32, 192, 32, 256, 32, 127>(
73 float* out, float* src0, float* src1, void* stream);74 float* out, float* src0, float* src1, void* stream);
74-template void LaunchTMULADDDST<float, 32, 128, 32, 192, 32, 256, 32, 127>(75+template void LaunchTMULA<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(
75- float* out, float* src0, float* src1, void* stream);
76-template void LaunchTMULADDDST<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(
77 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);76 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);
78-template void LaunchTMULADDDST<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>(77+template void LaunchTMULA<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>(
79 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);78 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);
80-template void LaunchTMULADDDST<aclFloat16, 1, 16384, 1, 16384, 1, 16384, 1, 16384>(79+template void LaunchTMULA<aclFloat16, 1, 16384, 1, 16384, 1, 16384, 1, 16384>(
81 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);80 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);
Mtests/npu/a2a3/src/st/testcase/CMakeLists.txt+1-1
@@ -228,7 +228,7 @@ tpushpop_dir_both
228tpushpop_dir_both_concurrent228tpushpop_dir_both_concurrent
229tpushpop_subtile229tpushpop_subtile
230tmadd230tmadd
231-tmuladddst231+tmula
232syncall232syncall
233)233)
234 234 
Rtests/cpu/st/testcase/tmuladddst/CMakeLists.txt→tests/npu/a2a3/src/st/testcase/tmula/CMakeLists.txt+1-1
@@ -8,4 +8,4 @@
8# See LICENSE in the root of the software repository for the full text of the License.8# See LICENSE in the root of the software repository for the full text of the License.
9# --------------------------------------------------------------------------------9# --------------------------------------------------------------------------------
10 10 
11-pto_cpu_sim_st(tmuladddst)11+pto_vec_st(tmula)
Rtests/npu/a2a3/src/st/testcase/tmuladddst/gen_data.py→tests/npu/a2a3/src/st/testcase/tmula/gen_data.py+1-1
@@ -78,7 +78,7 @@ class TestParams:
78 np.uint16: 'uint16',78 np.uint16: 'uint16',
79 np.uint8: 'uint8'79 np.uint8: 'uint8'
80 }[dtype]80 }[dtype]
81- self.name = f"TMULADDDSTTest.case_{dtype_str}_{dst_tile_row}x{dst_tile_col}_\81+ self.name = f"TMULATest.case_{dtype_str}_{dst_tile_row}x{dst_tile_col}_\
82{src0_tile_row}x{src0_tile_col}_{src1_tile_row}x{src1_tile_col}_\82{src0_tile_row}x{src0_tile_col}_{src1_tile_row}x{src1_tile_col}_\
83{valid_row}x{valid_col}"83{valid_row}x{valid_col}"
84 84 
Rtests/npu/a2a3/src/st/testcase/tmuladddst/main.cpp→tests/npu/a2a3/src/st/testcase/tmula/main.cpp+9-15
@@ -15,7 +15,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
15using namespace std;15using namespace std;
16using namespace PtoTestCommon;16using namespace PtoTestCommon;
17 17 
18-class TMULADDDSTTest : public testing::Test {18+class TMULATest : public testing::Test {
19protected:19protected:
20 void SetUp() override {}20 void SetUp() override {}
21 void TearDown() override {}21 void TearDown() override {}
@@ -33,12 +33,12 @@ std::string GetGoldenDir()
33template <33template <
34 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,34 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
35 int vCols, bool isHalf = true>35 int vCols, bool isHalf = true>
36-void LaunchTMULADDDST(T* out, T* src0, T* src1, void* stream);36+void LaunchTMULA(T* out, T* src0, T* src1, void* stream);
37 37 
38template <38template <
39 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,39 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
40 int vCols, bool isHalf = false>40 int vCols, bool isHalf = false>
41-void test_TMULADDDST()41+void test_TMULA()
42{42{
43 size_t fileSizeDst = dstTileH * dstTileW * sizeof(T);43 size_t fileSizeDst = dstTileH * dstTileW * sizeof(T);
44 size_t fileSizeSrc0 = src0TileH * src0TileW * sizeof(T);44 size_t fileSizeSrc0 = src0TileH * src0TileW * sizeof(T);
@@ -68,7 +68,7 @@ void test_TMULADDDST()
68 aclrtMemcpy(src0Device, fileSizeSrc0, src0Host, fileSizeSrc0, ACL_MEMCPY_HOST_TO_DEVICE);68 aclrtMemcpy(src0Device, fileSizeSrc0, src0Host, fileSizeSrc0, ACL_MEMCPY_HOST_TO_DEVICE);
69 aclrtMemcpy(src1Device, fileSizeSrc1, src1Host, fileSizeSrc1, ACL_MEMCPY_HOST_TO_DEVICE);69 aclrtMemcpy(src1Device, fileSizeSrc1, src1Host, fileSizeSrc1, ACL_MEMCPY_HOST_TO_DEVICE);
70 70 
71- LaunchTMULADDDST<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(71+ LaunchTMULA<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(
72 dstDevice, src0Device, src1Device, stream);72 dstDevice, src0Device, src1Device, stream);
73 73 
74 aclrtSynchronizeStream(stream);74 aclrtSynchronizeStream(stream);
@@ -97,16 +97,10 @@ void test_TMULADDDST()
97 EXPECT_TRUE(ret);97 EXPECT_TRUE(ret);
98}98}
99 99 
100-TEST_F(TMULADDDSTTest, case_float_64x64_64x64_64x64_64x64) { test_TMULADDDST<float, 64, 64, 64, 64, 64, 64, 64, 64>(); }100+TEST_F(TMULATest, case_float_64x64_64x64_64x64_64x64) { test_TMULA<float, 64, 64, 64, 64, 64, 64, 64, 64>(); }
101-TEST_F(TMULADDDSTTest, case_float_32x128_32x192_32x256_32x127)101+TEST_F(TMULATest, case_float_32x128_32x192_32x256_32x127) { test_TMULA<float, 32, 128, 32, 192, 32, 256, 32, 127>(); }
102+TEST_F(TMULATest, case_half_64x64_64x64_64x64_64x64) { test_TMULA<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(); }
103+TEST_F(TMULATest, case_half_32x128_32x192_32x256_32x127)
102{104{
103- test_TMULADDDST<float, 32, 128, 32, 192, 32, 256, 32, 127>();105+ test_TMULA<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>();
104-}
105-TEST_F(TMULADDDSTTest, case_half_64x64_64x64_64x64_64x64)
106-{
107- test_TMULADDDST<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>();
108-}
109-TEST_F(TMULADDDSTTest, case_half_32x128_32x192_32x256_32x127)
110-{
111- test_TMULADDDST<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>();
112}106}
Rtests/npu/a2a3/src/st/testcase/tmuladddst/tmuladddst_kernel.cpp→tests/npu/a2a3/src/st/testcase/tmula/tmula_kernel.cpp+9-10
@@ -17,7 +17,7 @@ using namespace pto;
17template <17template <
18 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,18 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
19 int vCols>19 int vCols>
20-__global__ AICORE void runTMULADDDST(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1)20+__global__ AICORE void runTMULA(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1)
21{21{
22 using DynShape = pto::Shape<-1, -1, -1, -1, -1>;22 using DynShape = pto::Shape<-1, -1, -1, -1, -1>;
23 using DynStride = pto::Stride<-1, -1, -1, -1, -1>;23 using DynStride = pto::Stride<-1, -1, -1, -1, -1>;
@@ -49,7 +49,7 @@ __global__ AICORE void runTMULADDDST(__gm__ T __out__* out, __gm__ T __in__* src
49 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);49 set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
50 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);50 wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
51#endif51#endif
52- TMULADDDST<TileDataDst, TileDataSrc0, TileDataSrc1>(dstTile, src0Tile, src1Tile);52+ TMULA<TileDataDst, TileDataSrc0, TileDataSrc1>(dstTile, src0Tile, src1Tile);
53#ifndef __PTO_AUTO__53#ifndef __PTO_AUTO__
54 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);54 set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
55 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);55 wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
@@ -61,22 +61,21 @@ __global__ AICORE void runTMULADDDST(__gm__ T __out__* out, __gm__ T __in__* src
61template <61template <
62 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,62 typename T, int dstTileH, int dstTileW, int src0TileH, int src0TileW, int src1TileH, int src1TileW, int vRows,
63 int vCols, bool isHalf = true>63 int vCols, bool isHalf = true>
64-void LaunchTMULADDDST(T* out, T* src0, T* src1, void* stream)64+void LaunchTMULA(T* out, T* src0, T* src1, void* stream)
65{65{
66 if constexpr (std::is_same_v<T, aclFloat16>) {66 if constexpr (std::is_same_v<T, aclFloat16>) {
67- runTMULADDDST<half, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>67+ runTMULA<half, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>
68 <<<1, nullptr, stream>>>((half*)(out), (half*)(src0), (half*)(src1));68 <<<1, nullptr, stream>>>((half*)(out), (half*)(src0), (half*)(src1));
69 } else {69 } else {
70- runTMULADDDST<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>70+ runTMULA<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>
71 <<<1, nullptr, stream>>>(out, src0, src1);71 <<<1, nullptr, stream>>>(out, src0, src1);
72 }72 }
73}73}
74 74 
75-template void LaunchTMULADDDST<float, 64, 64, 64, 64, 64, 64, 64, 64>(75+template void LaunchTMULA<float, 64, 64, 64, 64, 64, 64, 64, 64>(float* out, float* src0, float* src1, void* stream);
76+template void LaunchTMULA<float, 32, 128, 32, 192, 32, 256, 32, 127>(
76 float* out, float* src0, float* src1, void* stream);77 float* out, float* src0, float* src1, void* stream);
77-template void LaunchTMULADDDST<float, 32, 128, 32, 192, 32, 256, 32, 127>(78+template void LaunchTMULA<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(
78- float* out, float* src0, float* src1, void* stream);
79-template void LaunchTMULADDDST<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(
80 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);79 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);
81-template void LaunchTMULADDDST<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>(80+template void LaunchTMULA<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>(
82 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);81 aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);
Mtests/npu/a5/src/st/testcase/CMakeLists.txt+1-1
@@ -165,7 +165,7 @@ set(ALL_TESTCASES
165 tmov_mx165 tmov_mx
166 tmul166 tmul
167 tmadd167 tmadd
168- tmuladddst168+ tmula
169 tlrelu169 tlrelu
170 tdiv170 tdiv
171 tcmps171 tcmps
Atests/npu/a5/src/st/testcase/tmula+1-0
@@ -0,0 +1 @@
1+../../../../a2a3/src/st/testcase/tmula/
Dtests/npu/a5/src/st/testcase/tmuladddst+0-1
@@ -1 +0,0 @@
1-../../../../a2a3/src/st/testcase/tmuladddst/
Mtests/run_st.sh+1-1
@@ -531,7 +531,7 @@ if [ "$ENABLE_A3" = "true" ]; then # A2A3
531 python3 tests/script/run_st.py $ARGS -w -v a3 -t tmadd -g TMADDTest.case_float_32x128_32x192_32x256_32x127531 python3 tests/script/run_st.py $ARGS -w -v a3 -t tmadd -g TMADDTest.case_float_32x128_32x192_32x256_32x127
532 fi532 fi
533 if [ -z "$ST_PART" ] || [ "$ST_PART" = "2" ]; then533 if [ -z "$ST_PART" ] || [ "$ST_PART" = "2" ]; then
534- python3 tests/script/run_st.py $ARGS -w -v a3 -t tmuladddst -g TMULADDDSTTest.case_float_32x128_32x192_32x256_32x127534+ python3 tests/script/run_st.py $ARGS -w -v a3 -t tmula -g TMULATest.case_float_32x128_32x192_32x256_32x127
535 fi535 fi
536 if [ -z "$ST_PART" ] || [ "$ST_PART" = "2" ]; then536 if [ -z "$ST_PART" ] || [ "$ST_PART" = "2" ]; then
537 python3 tests/script/run_st.py $ARGS -w -v a3 -t textract_ndto2xnz -g TExtractNd2xNzTest.case_half537 python3 tests/script/run_st.py $ARGS -w -v a3 -t textract_ndto2xnz -g TExtractNd2xNzTest.case_half