revert: restore TMULADDDST ISA API-pto-isa-AtomGit

Mdocs/PTOISA.md+1-1

Mdocs/PTOISA_zh.md+1-1

Mdocs/coding/version-compatibility.md+0-1

	@@ -147,7 +147,6 @@ 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` \|
151		150
152	Update 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.	151	Update 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.
153		152

Mdocs/coding/version-compatibility_zh.md+0-1

	@@ -147,7 +147,6 @@ 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` \|
151		150
152	请同步更新 C++ intrinsic 调用、ISA 文档引用、生成图和 ST 用例名称。重命名后不再保留历史名称作为公开 wrapper。	151	请同步更新 C++ intrinsic 调用、ISA 文档引用、生成图和 ST 用例名称。重命名后不再保留历史名称作为公开 wrapper。
153		152

Rdocs/figures/isa/TMULA.svg→docs/figures/isa/TMULADDDST.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="TMULA 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="TMULADDDST 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">TMULA</text>	35	+<text x="40" y="46" class="title">TMULADDDST</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

Mdocs/isa/README_zh.md+2-2

Rdocs/isa/TMULA.md→docs/isa/TMULADDDST.md+93-95

	@@ -1,95 +1,93 @@
1	-# TMULA	1	+# TMULADDDST
2	-	2	+
3	-## Tile Operation Diagram	3	+## Tile Operation Diagram
4	-	4	+
5	-![TMULA tile operation](../figures/isa/TMULA.svg)	5	+![TMULADDDST tile operation](../figures/isa/TMULADDDST.svg)
6	-	6	+
7	-## Introduction	7	+## Introduction
8	-	8	+
9	-Elementwise operation: `src0 * src1 + dst`.	9	+Elementwise operation: `src0 * src1 + dst`.
10	-	10	+
11	-## Math Interpretation	11	+## 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 Syntax	17	+## Assembly Syntax
18	-	18	+
19	-Synchronous form:	19	+Synchronous form:
20	-	20	+
21	-```text	21	+```text
22	-%dst = tmula %src0, %src1 : !pto.tile<...>	22	+%dst = tmuladddst %src0, %src1 : !pto.tile<...>
23	-```	23	+```
24	-	24	+
25	-### AS Level 1 (SSA)	25	+### AS Level 1 (SSA)
26	-	26	+
27	-```text	27	+```text
28	-%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>	28	+%dst = pto.tmuladddst %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	-```text	33	+```text
34	-pto.tmula ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)	34	+pto.tmuladddst ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
35	-```	35	+```
36	-## C++ Intrinsic	36	+## 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	-```cpp	40	+```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 TMULA(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &...events);	42	+PTO_INST RecordEvent TMULADDDST(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &...events);
43	-```	43	+```
44	-	44	+
45	-## Constraints	45	+## 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	-- Common constraints:	50	+- Valid region:
51	- - Tile location must be vector (`TileData::Loc == TileType::Vec`).	51	+ - The op uses `dst.GetValidRow()` / `dst.GetValidCol()` as the iteration domain; `src0/src1` are assumed to be compatible (not validated by explicit runtime checks in this op).
52	- - Static valid bounds: `TileData::ValidRow <= TileData::Rows` and `TileData::ValidCol <= TileData::Cols`.	52	+- The op iterates over `dst.GetValidRow()` / `dst.GetValidCol()`.
53	- - Runtime valid shape: `dst`, `src0`, and `src1` must have the same valid row and column counts.	53	+
54	-- The op iterates over `dst.GetValidRow()` / `dst.GetValidCol()`.	54	+## Examples
55	-	55	+
56	-## Examples	56	+```cpp
57	-	57	+#include <pto/pto-inst.hpp>
58	-```cpp	58	+
59	-#include <pto/pto-inst.hpp>	59	+using namespace pto;
60	-	60	+
61	-using namespace pto;	61	+void example() {
62	-	62	+ using TileT = Tile<TileType::Vec, float, 16, 16>;
63	-void example() {	63	+ TileT a, b, out;
64	- using TileT = Tile<TileType::Vec, float, 16, 16>;	64	+ TMULADDDST(out, a, b);
65	- TileT a, b, out;	65	+}
66	- TMULA(out, a, b);	66	+```
67	-}	67	+
68	-```	68	+## ASM Form Examples
69	-	69	+
70	-## ASM Form Examples	70	+### Auto Mode
71	-	71	+
72	-### Auto Mode	72	+```text
73	-	73	+# Auto mode: compiler/runtime-managed placement and scheduling.
74	-```text	74	+%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
75	-# Auto mode: compiler/runtime-managed placement and scheduling.	75	+```
76	-%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>	76	+
77	-```	77	+### Manual Mode
78	-	78	+
79	-### Manual Mode	79	+```text
80	-	80	+# Manual mode: resources must be bound explicitly before issuing the instruction.
81	-```text	81	+# Optional for tile operands:
82	-# Manual mode: resources must be bound explicitly before issuing the instruction.	82	+# pto.tassign %arg0, @tile(0x1000)
83	-# Optional for tile operands:	83	+# pto.tassign %arg1, @tile(0x2000)
84	-# pto.tassign %arg0, @tile(0x1000)	84	+%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
85	-# pto.tassign %arg1, @tile(0x2000)	85	+```
86	-%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>	86	+
87	-```	87	+### PTO Assembly Form
88	-	88	+
89	-### PTO Assembly Form	89	+```text
90	-	90	+%dst = tmuladddst %src0, %src1 : !pto.tile<...>
91	-```text	91	+# AS Level 2 (DPS)
92	-%dst = tmula %src0, %src1 : !pto.tile<...>	92	+pto.tmuladddst ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
93	-# AS Level 2 (DPS)	93	+```
94	-pto.tmula ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
95	-```

Rdocs/isa/TMULA_zh.md→docs/isa/TMULADDDST_zh.md+98-97

	@@ -1,97 +1,98 @@
1	-# TMULA	1	+# TMULADDDST
2	-	2	+
3	-## 指令示意图	3	+## 指令示意图
4	-	4	+
5	-![TMULA tile operation](../figures/isa/TMULA.svg)	5	+![TMULADDDST tile operation](../figures/isa/TMULADDDST.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	-```text	21	+```text
22	-%dst = tmula %src0, %src1 : !pto.tile<...>, !pto.tile<...>	22	+%dst = tmuladddst %src0, %src1 : !pto.tile<...>, !pto.tile<...>
23	-```	23	+```
24	-	24	+
25	-### AS Level 1（SSA）	25	+### AS Level 1（SSA）
26	-	26	+
27	-```text	27	+```text
28	-%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>	28	+%dst = pto.tmuladddst %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	-```text	33	+```text
34	-pto.tmula ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)	34	+pto.tmuladddst 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	-```cpp	42	+```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 TMULA(TileDataDst &dst, TileDataSrc0 &src0, TileDataSrc1 &src1, WaitEvents &...events);	44	+PTO_INST RecordEvent TMULADDDST(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	-- 该操作在 `dst.GetValidRow()` / `dst.GetValidCol()` 上迭代。	56	+ - 标量类型必须与 Tile 数据类型一致。
57	-	57	+- 该操作在 `dst.GetValidRow()` / `dst.GetValidCol()` 上迭代。
58	-## 示例	58	+
59	-	59	+## 示例
60	-```cpp	60	+
61	-#include <pto/pto-inst.hpp>	61	+```cpp
62	-	62	+#include <pto/pto-inst.hpp>
63	-using namespace pto;	63	+
64	-	64	+using namespace pto;
65	-void example() {	65	+
66	- using TileT = Tile<TileType::Vec, float, 16, 16>;	66	+void example() {
67	- TileT a, b, out;	67	+ using TileT = Tile<TileType::Vec, float, 16, 16>;
68	- TMULA(out, a, b);	68	+ TileT a, b, out;
69	-}	69	+ TMULADDDST(out, a, b);
70	-```	70	+}
71	-	71	+```
72	-## 汇编示例（ASM）	72	+
73	-	73	+## 汇编示例（ASM）
74	-### 自动模式	74	+
75	-	75	+### 自动模式
76	-```text	76	+
77	-# 自动模式：由编译器/运行时负责资源放置与调度。	77	+```text
78	-%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>	78	+# 自动模式：由编译器/运行时负责资源放置与调度。
79	-```	79	+%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
80	-	80	+```
81	-### 手动模式	81	+
82	-	82	+### 手动模式
83	-```text	83	+
84	-# 手动模式：先显式绑定资源，再发射指令。	84	+```text
85	-# 可选（当该指令包含 tile 操作数时）：	85	+# 手动模式：先显式绑定资源，再发射指令。
86	-# pto.tassign %arg0, @tile(0x1000)	86	+# 可选（当该指令包含 tile 操作数时）：
87	-# pto.tassign %arg1, @tile(0x2000)	87	+# pto.tassign %arg0, @tile(0x1000)
88	-%dst = pto.tmula %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>	88	+# pto.tassign %arg1, @tile(0x2000)
89	-```	89	+%dst = pto.tmuladddst %src0, %src1 : (!pto.tile<...>, !pto.tile<...>) -> !pto.tile<...>
90	-	90	+```
91	-### PTO 汇编形式	91	+
92	-	92	+### PTO 汇编形式
93	-```text	93	+
94	-%dst = tmula %src0, %src1 : !pto.tile<...>, !pto.tile<...>	94	+```text
95	-# AS Level 2 (DPS)	95	+%dst = tmuladddst %src0, %src1 : !pto.tile<...>, !pto.tile<...>
96	-pto.tmula ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)	96	+# AS Level 2 (DPS)
97	-```	97	+pto.tmuladddst ins(%src0, %src1 : !pto.tile_buf<...>, !pto.tile_buf<...>) outs(%dst : !pto.tile_buf<...>)
		98	+```

Mdocs/isa/manifest.yaml+13-13

	@@ -172,19 +172,6 @@
172	],	172	],
173	"notes": []	173	"notes": []
174	},	174	},
175	- {
176	- "instruction": "TMULA",
177	- "category": "Elementwise (Tile-Tile)",
178	- "summary_en": "Elementwise ternary op: `src0 * src1 + dst`.",
179	- "summary_zh": "三元逐元素运算：`src0 * src1 + dst`。",
180	- "diagram_template": "elementwise",
181	- "operands": [
182	- "dst",
183	- "src0",
184	- "src1"
185	- ],
186	- "notes": []
187	- },
188	{	175	{
189	"instruction": "TMIN",	176	"instruction": "TMIN",
190	"category": "Elementwise (Tile-Tile)",	177	"category": "Elementwise (Tile-Tile)",
	@@ -459,6 +446,19 @@
459	],	446	],
460	"notes": []	447	"notes": []
461	},	448	},
		449	+ {
		450	+ "instruction": "TMULADDDST",
		451	+ "category": "Elementwise (Tile-Tile)",
		452	+ "summary_en": "Elementwise ternary op: `src0 * src1 + dst`.",
		453	+ "summary_zh": "三元逐元素运算：`src0 * src1 + dst`。",
		454	+ "diagram_template": "elementwise",
		455	+ "operands": [
		456	+ "dst",
		457	+ "src0",
		458	+ "src1"
		459	+ ],
		460	+ "notes": []
		461	+ },
462	{	462	{
463	"instruction": "TEXPANDS",	463	"instruction": "TEXPANDS",
464	"category": "Tile-Scalar / Tile-Immediate",	464	"category": "Tile-Scalar / Tile-Immediate",

Mdocs/mkdocs/gen_pages.py+1-1

	@@ -292,7 +292,6 @@ NAV_PAGES_EN = [
292	"docs/isa/TSUB.md",	292	"docs/isa/TSUB.md",
293	"docs/isa/TMUL.md",	293	"docs/isa/TMUL.md",
294	"docs/isa/TMADD.md",	294	"docs/isa/TMADD.md",
295	- "docs/isa/TMULA.md",
296	"docs/isa/TMIN.md",	295	"docs/isa/TMIN.md",
297	"docs/isa/TMAX.md",	296	"docs/isa/TMAX.md",
298	"docs/isa/TCMP.md",	297	"docs/isa/TCMP.md",
	@@ -303,6 +302,7 @@ NAV_PAGES_EN = [
303	"docs/isa/TLOG.md",	302	"docs/isa/TLOG.md",
304	"docs/isa/TRECIP.md",	303	"docs/isa/TRECIP.md",
305	"docs/isa/TPRELU.md",	304	"docs/isa/TPRELU.md",
		305	+ "docs/isa/TMULADDDST.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

	@@ -87,7 +87,6 @@ nav:
87	- TSUB: docs/isa/TSUB.md	87	- TSUB: docs/isa/TSUB.md
88	- TMUL: docs/isa/TMUL.md	88	- TMUL: docs/isa/TMUL.md
89	- TMADD: docs/isa/TMADD.md	89	- TMADD: docs/isa/TMADD.md
90	- - TMULA: docs/isa/TMULA.md
91	- TMIN: docs/isa/TMIN.md	90	- TMIN: docs/isa/TMIN.md
92	- TMAX: docs/isa/TMAX.md	91	- TMAX: docs/isa/TMAX.md
93	- TCMP: docs/isa/TCMP.md	92	- TCMP: docs/isa/TCMP.md
	@@ -98,6 +97,7 @@ nav:
98	- TLOG: docs/isa/TLOG.md	97	- TLOG: docs/isa/TLOG.md
99	- TRECIP: docs/isa/TRECIP.md	98	- TRECIP: docs/isa/TRECIP.md
100	- TPRELU: docs/isa/TPRELU.md	99	- TPRELU: docs/isa/TPRELU.md
		100	+ - TMULADDDST: docs/isa/TMULADDDST.md
101	- TCVT: docs/isa/TCVT.md	101	- TCVT: docs/isa/TCVT.md
102	- TSEL: docs/isa/TSEL.md	102	- TSEL: docs/isa/TSEL.md
103	- TRSQRT: docs/isa/TRSQRT.md	103	- TRSQRT: docs/isa/TRSQRT.md

Mdocs/mkdocs/src/manual/appendix-d-instruction-family-matrix.md+1-1

Mdocs/mkdocs/src/manual/appendix-d-instruction-family-matrix_zh.md+1-1

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`, and `TMULADDDST` to `TMULA`.")	208	+ lines.append("- Rename legacy fused multiply-add APIs: `TFUSEDMULADD` to `TMADD`.")
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`，`TMULADDDST` 改为 `TMULA`。")	248	+ lines.append("- 将历史融合乘加接口重命名：`TFUSEDMULADD` 改为 `TMADD`。")
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
630	TERNARY_ELEMENTWISE_SPEC = {	630	TERNARY_ELEMENTWISE_SPEC = {
631	"TMADD": ("src0 * dst(old) + src1", ["src0", "dst(old)", "src1"]),	631	"TMADD": ("src0 * dst(old) + src1", ["src0", "dst(old)", "src1"]),
632	- "TMULA": ("src0 * src1 + dst(old)", ["src0", "src1", "dst(old)"]),	632	+ "TMULADDDST": ("src0 * src1 + dst(old)", ["src0", "src1", "dst(old)"]),
633	}	633	}
634		634
635		635

Minclude/README.md+1-1

	@@ -102,7 +102,6 @@ This table tracks per-instruction backend availability:
102	\| [`TMOV_FP`](../docs/isa/TMOV_FP.md) \| Yes \| TODO \| Yes \| Yes \| Yes \| Yes \|	102	\| [`TMOV_FP`](../docs/isa/TMOV_FP.md) \| Yes \| TODO \| Yes \| Yes \| Yes \| Yes \|
103	\| [`TMRGSORT`](../docs/isa/TMRGSORT.md) \| Yes \| Yes \| Yes \| Yes \| Yes \| Yes \|	103	\| [`TMRGSORT`](../docs/isa/TMRGSORT.md) \| Yes \| Yes \| Yes \| Yes \| Yes \| Yes \|
104	\| [`TMUL`](../docs/isa/TMUL.md) \| Yes \| Yes \| Yes \| Yes \| Yes \| Yes \|	104	\| [`TMUL`](../docs/isa/TMUL.md) \| Yes \| Yes \| Yes \| Yes \| Yes \| Yes \|
105	-\| [`TMULA`](../docs/isa/TMULA.md) \| TODO \| TODO \| Yes \| Yes \| Yes \| TODO \|
106	\| [`TMULS`](../docs/isa/TMULS.md) \| Yes \| Yes \| Yes \| Yes \| Yes \| Yes \|	105	\| [`TMULS`](../docs/isa/TMULS.md) \| Yes \| Yes \| Yes \| Yes \| Yes \| Yes \|
107	\| [`TNEG`](../docs/isa/TNEG.md) \| Yes \| Yes \| Yes \| Yes \| Yes \| Yes \|	106	\| [`TNEG`](../docs/isa/TNEG.md) \| Yes \| Yes \| Yes \| Yes \| Yes \| Yes \|
108	\| [`TNOT`](../docs/isa/TNOT.md) \| Yes \| TODO \| Yes \| Yes \| Yes \| Yes \|	107	\| [`TNOT`](../docs/isa/TNOT.md) \| Yes \| TODO \| Yes \| Yes \| Yes \| Yes \|
	@@ -168,6 +167,7 @@ This table tracks per-instruction backend availability:
168	\| [`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 \|
169	\| [`TXOR`](../docs/isa/TXOR.md) \| Yes \| TODO \| Yes \| Yes \| Yes \| Yes \|	168	\| [`TXOR`](../docs/isa/TXOR.md) \| Yes \| TODO \| Yes \| Yes \| Yes \| Yes \|
170	\| [`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 \|
171		171
172	Notes:	172	Notes:
173		173

Minclude/README_zh.md+1-1

	@@ -102,7 +102,6 @@ PTO Tile Lib 对外的 C/C++ 头文件（以模板化、基本 header-only 为
102	\| [`TMOV_FP`](../docs/isa/TMOV_FP_zh.md) \| 是 \| TODO \| 是 \| 是 \| 是 \| 是 \|	102	\| [`TMOV_FP`](../docs/isa/TMOV_FP_zh.md) \| 是 \| TODO \| 是 \| 是 \| 是 \| 是 \|
103	\| [`TMRGSORT`](../docs/isa/TMRGSORT_zh.md) \| 是 \| 是 \| 是 \| 是 \| 是 \| 是 \|	103	\| [`TMRGSORT`](../docs/isa/TMRGSORT_zh.md) \| 是 \| 是 \| 是 \| 是 \| 是 \| 是 \|
104	\| [`TMUL`](../docs/isa/TMUL_zh.md) \| 是 \| 是 \| 是 \| 是 \| 是 \| 是 \|	104	\| [`TMUL`](../docs/isa/TMUL_zh.md) \| 是 \| 是 \| 是 \| 是 \| 是 \| 是 \|
105	-\| [`TMULA`](../docs/isa/TMULA_zh.md) \| TODO \| TODO \| 是 \| 是 \| 是 \| TODO \|
106	\| [`TMULS`](../docs/isa/TMULS_zh.md) \| 是 \| 是 \| 是 \| 是 \| 是 \| 是 \|	105	\| [`TMULS`](../docs/isa/TMULS_zh.md) \| 是 \| 是 \| 是 \| 是 \| 是 \| 是 \|
107	\| [`TNEG`](../docs/isa/TNEG_zh.md) \| 是 \| 是 \| 是 \| 是 \| 是 \| 是 \|	106	\| [`TNEG`](../docs/isa/TNEG_zh.md) \| 是 \| 是 \| 是 \| 是 \| 是 \| 是 \|
108	\| [`TNOT`](../docs/isa/TNOT_zh.md) \| 是 \| TODO \| 是 \| 是 \| 是 \| 是 \|	107	\| [`TNOT`](../docs/isa/TNOT_zh.md) \| 是 \| TODO \| 是 \| 是 \| 是 \| 是 \|
	@@ -167,6 +166,7 @@ PTO Tile Lib 对外的 C/C++ 头文件（以模板化、基本 header-only 为
167	\| [`TWAIT`](../docs/isa/comm/TWAIT_zh.md) \| 是 \| TODO \| 是 \| 是 \| 是 \| TODO \|	166	\| [`TWAIT`](../docs/isa/comm/TWAIT_zh.md) \| 是 \| TODO \| 是 \| 是 \| 是 \| TODO \|
168	\| [`TXOR`](../docs/isa/TXOR_zh.md) \| 是 \| TODO \| 是 \| 是 \| 是 \| 是 \|	167	\| [`TXOR`](../docs/isa/TXOR_zh.md) \| 是 \| TODO \| 是 \| 是 \| 是 \| 是 \|
169	\| [`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 \|
170		170
171	说明：	171	说明：
172		172

Minclude/pto/common/event.hpp+2-2

	@@ -30,7 +30,6 @@ enum class Op : uint16_t {
30	TSUB,	30	TSUB,
31	TMUL,	31	TMUL,
32	TMADD,	32	TMADD,
33	- TMULA,
34	TMULS,	33	TMULS,
35	TDIV,	34	TDIV,
36	TDIVS,	35	TDIVS,
	@@ -138,6 +137,7 @@ enum class Op : uint16_t {
138	TCOLARGMIN,	137	TCOLARGMIN,
139	TROWARGMAX,	138	TROWARGMAX,
140	TROWARGMIN,	139	TROWARGMIN,
		140	+ TMULADDDST,
141	TPREFETCH,	141	TPREFETCH,
142	TFILLPAD_VEC,	142	TFILLPAD_VEC,
143	TFILLPAD_MAT,	143	TFILLPAD_MAT,
	@@ -173,7 +173,6 @@ PTO_DEFINE_OP_PIPE(Op::TAXPY, PIPE_V);
173	PTO_DEFINE_OP_PIPE(Op::TSUB, PIPE_V);	173	PTO_DEFINE_OP_PIPE(Op::TSUB, PIPE_V);
174	PTO_DEFINE_OP_PIPE(Op::TMUL, PIPE_V);	174	PTO_DEFINE_OP_PIPE(Op::TMUL, PIPE_V);
175	PTO_DEFINE_OP_PIPE(Op::TMADD, PIPE_V);	175	PTO_DEFINE_OP_PIPE(Op::TMADD, PIPE_V);
176	-PTO_DEFINE_OP_PIPE(Op::TMULA, PIPE_V);
177	PTO_DEFINE_OP_PIPE(Op::TMULS, PIPE_V);	176	PTO_DEFINE_OP_PIPE(Op::TMULS, PIPE_V);
178	PTO_DEFINE_OP_PIPE(Op::TDIV, PIPE_V);	177	PTO_DEFINE_OP_PIPE(Op::TDIV, PIPE_V);
179	PTO_DEFINE_OP_PIPE(Op::TDIVS, PIPE_V);	178	PTO_DEFINE_OP_PIPE(Op::TDIVS, PIPE_V);
	@@ -287,6 +286,7 @@ PTO_DEFINE_OP_PIPE(Op::TFILLPAD_MAT, PIPE_MTE2);
287	PTO_DEFINE_OP_PIPE(Op::MGATHER_VEC, PIPE_V);	286	PTO_DEFINE_OP_PIPE(Op::MGATHER_VEC, PIPE_V);
288	PTO_DEFINE_OP_PIPE(Op::MGATHER_MAT, PIPE_MTE2);	287	PTO_DEFINE_OP_PIPE(Op::MGATHER_MAT, PIPE_MTE2);
289	PTO_DEFINE_OP_PIPE(Op::MSCATTER, PIPE_V);	288	PTO_DEFINE_OP_PIPE(Op::MSCATTER, PIPE_V);
		289	+PTO_DEFINE_OP_PIPE(Op::TMULADDDST, PIPE_V);
290	PTO_DEFINE_OP_PIPE(Op::TQUANT, PIPE_V);	290	PTO_DEFINE_OP_PIPE(Op::TQUANT, PIPE_V);
291	PTO_DEFINE_OP_PIPE(Op::THISTOGRAM, PIPE_V);	291	PTO_DEFINE_OP_PIPE(Op::THISTOGRAM, PIPE_V);
292	PTO_DEFINE_OP_PIPE(Op::TINTERLEAVE, PIPE_V);	292	PTO_DEFINE_OP_PIPE(Op::TINTERLEAVE, PIPE_V);

Minclude/pto/common/pto_instr.hpp+8-8

	@@ -171,14 +171,6 @@ PTO_INST RecordEvent TMADD(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& s
171	return {};	171	return {};
172	}	172	}
173		173
174	-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>
175	-PTO_INST RecordEvent TMULA(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1, WaitEvents&... events)
176	-{
177	- detail::PtoWaitEvents(events...);
178	- MAP_INSTR_IMPL(TMULA, dst, src0, src1);
179	- return {};
180	-}
181	-
182	template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>	174	template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>
183	PTO_INST RecordEvent TMIN(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1, WaitEvents&... events)	175	PTO_INST RecordEvent TMIN(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1, WaitEvents&... events)
184	{	176	{
	@@ -1266,6 +1258,14 @@ PTO_INST RecordEvent TPARTARGMIN(
1266	return {};	1258	return {};
1267	}	1259	}
1268		1260
		1261	+template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>
		1262	+PTO_INST RecordEvent TMULADDDST(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1, WaitEvents&... events)
		1263	+{
		1264	+ detail::PtoWaitEvents(events...);
		1265	+ MAP_INSTR_IMPL(TMULADDDST, dst, src0, src1);
		1266	+ return {};
		1267	+}
		1268	+
1269	template <typename TileDataD, typename TileDataS, typename TmpTileData>	1269	template <typename TileDataD, typename TileDataS, typename TmpTileData>
1270	PTO_INTERNAL void TCVT_IMPL(	1270	PTO_INTERNAL void TCVT_IMPL(
1271	TileDataD& dst, TileDataS& src, TmpTileData& tmp, RoundMode mode, SaturationMode satMode, bool needSetCtrl);	1271	TileDataD& dst, TileDataS& src, TmpTileData& tmp, RoundMode mode, SaturationMode satMode, bool needSetCtrl);

Minclude/pto/common/pto_instr_impl.hpp+2-3

Minclude/pto/costmodel/pto_instr.hpp+0-8

	@@ -365,14 +365,6 @@ PTO_INST RecordEvent TMADD(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& s
365	return {};	365	return {};
366	}	366	}
367		367
368	-template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>
369	-PTO_INST RecordEvent TMULA(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1, WaitEvents&... events)
370	-{
371	- detail::PtoWaitEvents(events...);
372	- MAP_INSTR_IMPL(TMULA, dst, src0, src1);
373	- return {};
374	-}
375	-
376	template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>	368	template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1, typename... WaitEvents>
377	PTO_INST RecordEvent TMIN(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1, WaitEvents&... events)	369	PTO_INST RecordEvent TMIN(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1, WaitEvents&... events)
378	{	370	{

Minclude/pto/cpu/ElementOp.h+2-2

	@@ -25,7 +25,7 @@ enum class ElementOp {
25	OP_SUB,	25	OP_SUB,
26	OP_MUL,	26	OP_MUL,
27	OP_DIV,	27	OP_DIV,
28	- OP_MULA,	28	+ OP_MULADDDST,
29	OP_MADD,	29	OP_MADD,
30	OP_REM,	30	OP_REM,
31	OP_SHL,	31	OP_SHL,
	@@ -121,7 +121,7 @@ struct ElementOpCal<DType, ElementOp::OP_DIV> {
121	};	121	};
122		122
123	template <typename DType>	123	template <typename DType>
124	-struct ElementOpCal<DType, ElementOp::OP_MULA> {	124	+struct ElementOpCal<DType, ElementOp::OP_MULADDDST> {
125	static void apply(DType& dst, DType& src0, DType& src1, size_t) { dst = static_cast<DType>(src0 * src1) + dst; }	125	static void apply(DType& dst, DType& src0, DType& src1, size_t) { dst = static_cast<DType>(src0 * src1) + dst; }
126		126
127	static void apply(DType& dst, const DType& src0, const DType& src1) { dst = static_cast<DType>(src0 * src1) + dst; }	127	static void apply(DType& dst, const DType& src0, const DType& src1) { dst = static_cast<DType>(src0 * src1) + dst; }

Minclude/pto/cpu/ElementTileOp.h+1-1

	@@ -114,7 +114,7 @@ BINARY_OP_DEF(AND)
114	BINARY_OP_DEF(OR)	114	BINARY_OP_DEF(OR)
115	BINARY_OP_DEF(XOR)	115	BINARY_OP_DEF(XOR)
116	BINARY_OP_DEF(MIN)	116	BINARY_OP_DEF(MIN)
117	-BINARY_OP_DEF(MULA)	117	+BINARY_OP_DEF(MULADDDST)
118	BINARY_OP_DEF(MADD)	118	BINARY_OP_DEF(MADD)
119		119
120	UNARY_OP_DEF(NEG)	120	UNARY_OP_DEF(NEG)

Rinclude/pto/npu/a2a3/TMula.hpp→include/pto/npu/a2a3/TMulAddDst.hpp+14-14

Rinclude/pto/npu/a5/TMula.hpp→include/pto/npu/a5/TMulAddDst.hpp+13-13

Mtests/cpu/st/testcase/CMakeLists.txt+1-1

	@@ -98,7 +98,7 @@ tmins
98	tmov	98	tmov
99	tmrgsort	99	tmrgsort
100	tmul	100	tmul
101	-tmula	101	+tmuladddst
102	tmuls	102	tmuls
103	tneg	103	tneg
104	tnot	104	tnot

Rtests/npu/a2a3/src/st/testcase/tmula/CMakeLists.txt→tests/cpu/st/testcase/tmuladddst/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(tmula)	11	+pto_cpu_sim_st(tmuladddst)

Rtests/cpu/st/testcase/tmula/gen_data.py→tests/cpu/st/testcase/tmuladddst/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"TMULATest.case_{dtype_str}_{dst_tile_row}x{dst_tile_col}_\	93	+ self.name = f"TMULADDDSTTest.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/tmula/main.cpp→tests/cpu/st/testcase/tmuladddst/main.cpp+17-11

	@@ -15,7 +15,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
15	using namespace std;	15	using namespace std;
16	using namespace PtoTestCommon;	16	using namespace PtoTestCommon;
17		17
18	-class TMULATest : public testing::Test {	18	+class TMULADDDSTTest : public testing::Test {
19	protected:	19	protected:
20	void SetUp() override {}	20	void SetUp() override {}
21	void TearDown() override {}	21	void TearDown() override {}
	@@ -33,12 +33,12 @@ std::string GetGoldenDir()
33	template <	33	template <
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 LaunchTMULA(T* out, T* src0, T* src1, void* stream);	36	+void LaunchTMULADDDST(T* out, T* src0, T* src1, void* stream);
37		37
38	template <	38	template <
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_TMULA()	41	+void test_TMULADDDST()
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_TMULA()
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	- LaunchTMULA<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(	71	+ LaunchTMULADDDST<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,15 +97,21 @@ void test_TMULA()
97	EXPECT_TRUE(ret);	97	EXPECT_TRUE(ret);
98	}	98	}
99		99
100	-TEST_F(TMULATest, case_float_64x64_64x64_64x64_64x64) { test_TMULA<float, 64, 64, 64, 64, 64, 64, 64, 64>(); }	100	+TEST_F(TMULADDDSTTest, case_float_64x64_64x64_64x64_64x64) { test_TMULADDDST<float, 64, 64, 64, 64, 64, 64, 64, 64>(); }
101	-TEST_F(TMULATest, case_float_32x128_32x192_32x256_32x127) { test_TMULA<float, 32, 128, 32, 192, 32, 256, 32, 127>(); }	101	+TEST_F(TMULADDDSTTest, case_float_32x128_32x192_32x256_32x127)
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)
104	{	102	{
105	- test_TMULA<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>();	103	+ test_TMULADDDST<float, 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>();
106	}	112	}
107		113
108	-TEST_F(TMULATest, case_half_1x16384_1x16384_1x16384_1x16384)	114	+TEST_F(TMULADDDSTTest, case_half_1x16384_1x16384_1x16384_1x16384)
109	{	115	{
110	- test_TMULA<aclFloat16, 1, 16384, 1, 16384, 1, 16384, 1, 16384>();	116	+ test_TMULADDDST<aclFloat16, 1, 16384, 1, 16384, 1, 16384, 1, 16384>();
111	}	117	}

Rtests/cpu/st/testcase/tmula/tmula_kernel.cpp→tests/cpu/st/testcase/tmuladddst/tmuladddst_kernel.cpp+11-10

Mtests/npu/a2a3/src/st/testcase/CMakeLists.txt+1-1

	@@ -228,7 +228,7 @@ tpushpop_dir_both
228	tpushpop_dir_both_concurrent	228	tpushpop_dir_both_concurrent
229	tpushpop_subtile	229	tpushpop_subtile
230	tmadd	230	tmadd
231	-tmula	231	+tmuladddst
232	syncall	232	syncall
233	)	233	)
234		234

Rtests/cpu/st/testcase/tmula/CMakeLists.txt→tests/npu/a2a3/src/st/testcase/tmuladddst/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(tmula)	11	+pto_vec_st(tmuladddst)

Rtests/npu/a2a3/src/st/testcase/tmula/gen_data.py→tests/npu/a2a3/src/st/testcase/tmuladddst/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"TMULATest.case_{dtype_str}_{dst_tile_row}x{dst_tile_col}_\	81	+ self.name = f"TMULADDDSTTest.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/tmula/main.cpp→tests/npu/a2a3/src/st/testcase/tmuladddst/main.cpp+15-9

	@@ -15,7 +15,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
15	using namespace std;	15	using namespace std;
16	using namespace PtoTestCommon;	16	using namespace PtoTestCommon;
17		17
18	-class TMULATest : public testing::Test {	18	+class TMULADDDSTTest : public testing::Test {
19	protected:	19	protected:
20	void SetUp() override {}	20	void SetUp() override {}
21	void TearDown() override {}	21	void TearDown() override {}
	@@ -33,12 +33,12 @@ std::string GetGoldenDir()
33	template <	33	template <
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 LaunchTMULA(T* out, T* src0, T* src1, void* stream);	36	+void LaunchTMULADDDST(T* out, T* src0, T* src1, void* stream);
37		37
38	template <	38	template <
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_TMULA()	41	+void test_TMULADDDST()
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_TMULA()
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	- LaunchTMULA<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(	71	+ LaunchTMULADDDST<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,10 +97,16 @@ void test_TMULA()
97	EXPECT_TRUE(ret);	97	EXPECT_TRUE(ret);
98	}	98	}
99		99
100	-TEST_F(TMULATest, case_float_64x64_64x64_64x64_64x64) { test_TMULA<float, 64, 64, 64, 64, 64, 64, 64, 64>(); }	100	+TEST_F(TMULADDDSTTest, case_float_64x64_64x64_64x64_64x64) { test_TMULADDDST<float, 64, 64, 64, 64, 64, 64, 64, 64>(); }
101	-TEST_F(TMULATest, case_float_32x128_32x192_32x256_32x127) { test_TMULA<float, 32, 128, 32, 192, 32, 256, 32, 127>(); }	101	+TEST_F(TMULADDDSTTest, case_float_32x128_32x192_32x256_32x127)
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)
104	{	102	{
105	- test_TMULA<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>();	103	+ test_TMULADDDST<float, 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>();
106	}	112	}

Rtests/npu/a2a3/src/st/testcase/tmula/tmula_kernel.cpp→tests/npu/a2a3/src/st/testcase/tmuladddst/tmuladddst_kernel.cpp+10-9

Mtests/npu/a5/src/st/testcase/CMakeLists.txt+1-1

	@@ -165,7 +165,7 @@ set(ALL_TESTCASES
165	tmov_mx	165	tmov_mx
166	tmul	166	tmul
167	tmadd	167	tmadd
168	- tmula	168	+ tmuladddst
169	tlrelu	169	tlrelu
170	tdiv	170	tdiv
171	tcmps	171	tcmps

Dtests/npu/a5/src/st/testcase/tmula+0-1

	@@ -1 +0,0 @@
1	-../../../../a2a3/src/st/testcase/tmula/

Atests/npu/a5/src/st/testcase/tmuladddst+1-0

	@@ -0,0 +1 @@
		1	+../../../../a2a3/src/st/testcase/tmuladddst/

Mtests/run_st.sh+1-4

	@@ -32,7 +32,6 @@ This page is the source-synchronized ISA index generated from `docs/isa/manifest
32	\| Elementwise (Tile-Tile) \| [`TSUB`](isa/TSUB.md) \| Elementwise subtract of two tiles. \|	32	\| Elementwise (Tile-Tile) \| [`TSUB`](isa/TSUB.md) \| Elementwise subtract of two tiles. \|
33	\| Elementwise (Tile-Tile) \| [`TMUL`](isa/TMUL.md) \| Elementwise multiply of two tiles. \|	33	\| Elementwise (Tile-Tile) \| [`TMUL`](isa/TMUL.md) \| Elementwise multiply of two tiles. \|
34	\| Elementwise (Tile-Tile) \| [`TMADD`](isa/TMADD.md) \| Elementwise ternary op: `src0 * dst + src1`. \|	34	\| Elementwise (Tile-Tile) \| [`TMADD`](isa/TMADD.md) \| Elementwise ternary op: `src0 * dst + src1`. \|
35	-\| Elementwise (Tile-Tile) \| [`TMULA`](isa/TMULA.md) \| Elementwise ternary op: `src0 * src1 + dst`. \|
36	\| Elementwise (Tile-Tile) \| [`TMIN`](isa/TMIN.md) \| Elementwise minimum of two tiles. \|	35	\| Elementwise (Tile-Tile) \| [`TMIN`](isa/TMIN.md) \| Elementwise minimum of two tiles. \|
37	\| Elementwise (Tile-Tile) \| [`TMAX`](isa/TMAX.md) \| Elementwise maximum of two tiles. \|	36	\| Elementwise (Tile-Tile) \| [`TMAX`](isa/TMAX.md) \| Elementwise maximum of two tiles. \|
38	\| Elementwise (Tile-Tile) \| [`TCMP`](isa/TCMP.md) \| Compare two tiles and write a packed predicate mask. \|	37	\| Elementwise (Tile-Tile) \| [`TCMP`](isa/TCMP.md) \| Compare two tiles and write a packed predicate mask. \|
	@@ -54,6 +53,7 @@ This page is the source-synchronized ISA index generated from `docs/isa/manifest
54	\| 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. \|
55	\| 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. \|
56	\| 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`. \|
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). \|

	@@ -32,7 +32,6 @@
32	\| 逐元素（Tile-Tile） \| [`TSUB`](isa/TSUB_zh.md) \| 两个 Tile 的逐元素减法。 \|	32	\| 逐元素（Tile-Tile） \| [`TSUB`](isa/TSUB_zh.md) \| 两个 Tile 的逐元素减法。 \|
33	\| 逐元素（Tile-Tile） \| [`TMUL`](isa/TMUL_zh.md) \| 两个 Tile 的逐元素乘法。 \|	33	\| 逐元素（Tile-Tile） \| [`TMUL`](isa/TMUL_zh.md) \| 两个 Tile 的逐元素乘法。 \|
34	\| 逐元素（Tile-Tile） \| [`TMADD`](isa/TMADD_zh.md) \| 三元逐元素运算：`src0 * dst + src1`。 \|	34	\| 逐元素（Tile-Tile） \| [`TMADD`](isa/TMADD_zh.md) \| 三元逐元素运算：`src0 * dst + src1`。 \|
35	-\| 逐元素（Tile-Tile） \| [`TMULA`](isa/TMULA_zh.md) \| 三元逐元素运算：`src0 * src1 + dst`。 \|
36	\| 逐元素（Tile-Tile） \| [`TMIN`](isa/TMIN_zh.md) \| 两个 Tile 的逐元素最小值。 \|	35	\| 逐元素（Tile-Tile） \| [`TMIN`](isa/TMIN_zh.md) \| 两个 Tile 的逐元素最小值。 \|
37	\| 逐元素（Tile-Tile） \| [`TMAX`](isa/TMAX_zh.md) \| 两个 Tile 的逐元素最大值。 \|	36	\| 逐元素（Tile-Tile） \| [`TMAX`](isa/TMAX_zh.md) \| 两个 Tile 的逐元素最大值。 \|
38	\| 逐元素（Tile-Tile） \| [`TCMP`](isa/TCMP_zh.md) \| 比较两个 Tile 并写入一个打包的谓词掩码。 \|	37	\| 逐元素（Tile-Tile） \| [`TCMP`](isa/TCMP_zh.md) \| 比较两个 Tile 并写入一个打包的谓词掩码。 \|
	@@ -54,6 +53,7 @@
54	\| 逐元素（Tile-Tile） \| [`TNEG`](isa/TNEG_zh.md) \| Tile 的逐元素取负。 \|	53	\| 逐元素（Tile-Tile） \| [`TNEG`](isa/TNEG_zh.md) \| Tile 的逐元素取负。 \|
55	\| 逐元素（Tile-Tile） \| [`TREM`](isa/TREM_zh.md) \| 两个 Tile 的逐元素余数，余数符号与除数相同。 \|	54	\| 逐元素（Tile-Tile） \| [`TREM`](isa/TREM_zh.md) \| 两个 Tile 的逐元素余数，余数符号与除数相同。 \|
56	\| 逐元素（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`。 \|
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 逐元素选择）。 \|

	@@ -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`, and `TMULADDDST` to `TMULA`.	36	+- Rename legacy fused multiply-add APIs: `TFUSEDMULADD` to `TMADD`.
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.
	@@ -57,7 +57,6 @@ Migration guidance:
57	- [TSUB](TSUB.md) - Elementwise subtract of two tiles.	57	- [TSUB](TSUB.md) - Elementwise subtract of two tiles.
58	- [TMUL](TMUL.md) - Elementwise multiply of two tiles.	58	- [TMUL](TMUL.md) - Elementwise multiply of two tiles.
59	- [TMADD](TMADD.md) - Elementwise ternary op: `src0 * dst + src1`.	59	- [TMADD](TMADD.md) - Elementwise ternary op: `src0 * dst + src1`.
60	-- [TMULA](TMULA.md) - Elementwise ternary op: `src0 * src1 + dst`.
61	- [TMIN](TMIN.md) - Elementwise minimum of two tiles.	60	- [TMIN](TMIN.md) - Elementwise minimum of two tiles.
62	- [TMAX](TMAX.md) - Elementwise maximum of two tiles.	61	- [TMAX](TMAX.md) - Elementwise maximum of two tiles.
63	- [TCMP](TCMP.md) - Compare two tiles and write a packed predicate mask.	62	- [TCMP](TCMP.md) - Compare two tiles and write a packed predicate mask.
	@@ -79,6 +78,7 @@ Migration guidance:
79	- [TNEG](TNEG.md) - Elementwise negation of a tile.	78	- [TNEG](TNEG.md) - Elementwise negation of a tile.
80	- [TREM](TREM.md) - Elementwise remainder of two tiles.	79	- [TREM](TREM.md) - Elementwise remainder of two tiles.
81	- [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`.
82		82
83	## Tile-Scalar / Tile-Immediate	83	## 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.

	@@ -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`，`TMULADDDST` 改为 `TMULA`。	36	+- 将历史融合乘加接口重命名：`TFUSEDMULADD` 改为 `TMADD`。
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 地址。
	@@ -57,7 +57,6 @@
57	- [TSUB](TSUB_zh.md) - 两个 Tile 的逐元素减法。	57	- [TSUB](TSUB_zh.md) - 两个 Tile 的逐元素减法。
58	- [TMUL](TMUL_zh.md) - 两个 Tile 的逐元素乘法。	58	- [TMUL](TMUL_zh.md) - 两个 Tile 的逐元素乘法。
59	- [TMADD](TMADD_zh.md) - 三元逐元素运算：`src0 * dst + src1`。	59	- [TMADD](TMADD_zh.md) - 三元逐元素运算：`src0 * dst + src1`。
60	-- [TMULA](TMULA_zh.md) - 三元逐元素运算：`src0 * src1 + dst`。
61	- [TMIN](TMIN_zh.md) - 两个 Tile 的逐元素最小值。	60	- [TMIN](TMIN_zh.md) - 两个 Tile 的逐元素最小值。
62	- [TMAX](TMAX_zh.md) - 两个 Tile 的逐元素最大值。	61	- [TMAX](TMAX_zh.md) - 两个 Tile 的逐元素最大值。
63	- [TCMP](TCMP_zh.md) - 比较两个 Tile 并写入一个打包的谓词掩码。	62	- [TCMP](TCMP_zh.md) - 比较两个 Tile 并写入一个打包的谓词掩码。
	@@ -79,6 +78,7 @@
79	- [TNEG](TNEG_zh.md) - Tile 的逐元素取负。	78	- [TNEG](TNEG_zh.md) - Tile 的逐元素取负。
80	- [TREM](TREM_zh.md) - 两个 Tile 的逐元素余数，余数符号与除数相同。	79	- [TREM](TREM_zh.md) - 两个 Tile 的逐元素余数，余数符号与除数相同。
81	- [TFMOD](TFMOD_zh.md) - 两个 Tile 的逐元素余数，余数符号与被除数相同。	80	- [TFMOD](TFMOD_zh.md) - 两个 Tile 的逐元素余数，余数符号与被除数相同。
		81	+- [TMULADDDST](TMULADDDST_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 中。

	@@ -45,7 +45,6 @@ This appendix is generated from `docs/isa/manifest.yaml` and provides a source-s
45	\| Elementwise (Tile-Tile) \| `TSUB` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TSUB.md` \|	45	\| Elementwise (Tile-Tile) \| `TSUB` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TSUB.md` \|
46	\| Elementwise (Tile-Tile) \| `TMUL` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMUL.md` \|	46	\| Elementwise (Tile-Tile) \| `TMUL` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMUL.md` \|
47	\| Elementwise (Tile-Tile) \| `TMADD` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMADD.md` \|	47	\| Elementwise (Tile-Tile) \| `TMADD` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMADD.md` \|
48	-\| Elementwise (Tile-Tile) \| `TMULA` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMULA.md` \|
49	\| Elementwise (Tile-Tile) \| `TMIN` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMIN.md` \|	48	\| Elementwise (Tile-Tile) \| `TMIN` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMIN.md` \|
50	\| Elementwise (Tile-Tile) \| `TMAX` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMAX.md` \|	49	\| Elementwise (Tile-Tile) \| `TMAX` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMAX.md` \|
51	\| Elementwise (Tile-Tile) \| `TCMP` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TCMP.md` \|	50	\| Elementwise (Tile-Tile) \| `TCMP` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TCMP.md` \|
	@@ -67,6 +66,7 @@ This appendix is generated from `docs/isa/manifest.yaml` and provides a source-s
67	\| 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` \|
68	\| 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` \|
69	\| 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` \|
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` \|

	@@ -45,7 +45,6 @@
45	\| 逐元素（Tile-Tile） \| `TSUB` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TSUB_zh.md` \|	45	\| 逐元素（Tile-Tile） \| `TSUB` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TSUB_zh.md` \|
46	\| 逐元素（Tile-Tile） \| `TMUL` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMUL_zh.md` \|	46	\| 逐元素（Tile-Tile） \| `TMUL` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMUL_zh.md` \|
47	\| 逐元素（Tile-Tile） \| `TMADD` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMADD_zh.md` \|	47	\| 逐元素（Tile-Tile） \| `TMADD` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMADD_zh.md` \|
48	-\| 逐元素（Tile-Tile） \| `TMULA` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMULA_zh.md` \|
49	\| 逐元素（Tile-Tile） \| `TMIN` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMIN_zh.md` \|	48	\| 逐元素（Tile-Tile） \| `TMIN` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMIN_zh.md` \|
50	\| 逐元素（Tile-Tile） \| `TMAX` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMAX_zh.md` \|	49	\| 逐元素（Tile-Tile） \| `TMAX` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TMAX_zh.md` \|
51	\| 逐元素（Tile-Tile） \| `TCMP` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TCMP_zh.md` \|	50	\| 逐元素（Tile-Tile） \| `TCMP` \| `elementwise` \| `dst, src0, src1` \| `docs/isa/TCMP_zh.md` \|
	@@ -67,6 +66,7 @@
67	\| 逐元素（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` \|
68	\| 逐元素（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` \|
69	\| 逐元素（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` \|
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` \|

	@@ -25,7 +25,6 @@ See LICENSE in the root of the software repository for the full text of the Lice
25	#include "pto/npu/a2a3/TSubS.hpp"	25	#include "pto/npu/a2a3/TSubS.hpp"
26	#include "pto/npu/a2a3/TMul.hpp"	26	#include "pto/npu/a2a3/TMul.hpp"
27	#include "pto/npu/a2a3/TMAdd.hpp"	27	#include "pto/npu/a2a3/TMAdd.hpp"
28	-#include "pto/npu/a2a3/TMula.hpp"
29	#include "pto/npu/a2a3/TMulS.hpp"	28	#include "pto/npu/a2a3/TMulS.hpp"
30	#include "pto/npu/a2a3/TMin.hpp"	29	#include "pto/npu/a2a3/TMin.hpp"
31	#include "pto/npu/a2a3/TMins.hpp"	30	#include "pto/npu/a2a3/TMins.hpp"
	@@ -125,7 +124,6 @@ See LICENSE in the root of the software repository for the full text of the Lice
125	#include "pto/npu/a2a3/TMov.hpp"	124	#include "pto/npu/a2a3/TMov.hpp"
126	#include "pto/npu/a2a3/TMul.hpp"	125	#include "pto/npu/a2a3/TMul.hpp"
127	#include "pto/npu/a2a3/TMAdd.hpp"	126	#include "pto/npu/a2a3/TMAdd.hpp"
128	-#include "pto/npu/a2a3/TMula.hpp"
129	#include "pto/npu/a2a3/TSort32.hpp"	127	#include "pto/npu/a2a3/TSort32.hpp"
130	#include "pto/npu/a2a3/TSel.hpp"	128	#include "pto/npu/a2a3/TSel.hpp"
131	#include "pto/npu/a2a3/TGather.hpp"	129	#include "pto/npu/a2a3/TGather.hpp"
	@@ -185,6 +183,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
185	#include "pto/npu/a2a3/TAlloc.hpp"	183	#include "pto/npu/a2a3/TAlloc.hpp"
186	#include "pto/npu/a2a3/TFree.hpp"	184	#include "pto/npu/a2a3/TFree.hpp"
187	#include "pto/npu/a2a3/TColReduceIdx.hpp"	185	#include "pto/npu/a2a3/TColReduceIdx.hpp"
		186	+#include "pto/npu/a2a3/TMulAddDst.hpp"
188	#endif	187	#endif
189	#endif	188	#endif
190		189
	@@ -290,7 +289,6 @@ See LICENSE in the root of the software repository for the full text of the Lice
290	#include "pto/npu/a5/TDiv.hpp"	289	#include "pto/npu/a5/TDiv.hpp"
291	#include "pto/npu/a5/TMul.hpp"	290	#include "pto/npu/a5/TMul.hpp"
292	#include "pto/npu/a5/TMAdd.hpp"	291	#include "pto/npu/a5/TMAdd.hpp"
293	-#include "pto/npu/a5/TMula.hpp"
294	#include "pto/npu/a5/TScatter.hpp"	292	#include "pto/npu/a5/TScatter.hpp"
295	#include "pto/npu/a5/MGather.hpp"	293	#include "pto/npu/a5/MGather.hpp"
296	#include "pto/npu/a5/MScatter.hpp"	294	#include "pto/npu/a5/MScatter.hpp"
	@@ -311,6 +309,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
311	#include "pto/npu/a5/TColReduceIdx.hpp"	309	#include "pto/npu/a5/TColReduceIdx.hpp"
312	#include "pto/npu/a5/TInterleave.hpp"	310	#include "pto/npu/a5/TInterleave.hpp"
313	#include "pto/npu/a5/TDeInterleave.hpp"	311	#include "pto/npu/a5/TDeInterleave.hpp"
		312	+#include "pto/npu/a5/TMulAddDst.hpp"
314	#endif // __COSTMODEL	313	#endif // __COSTMODEL
315	#endif	314	#endif
316		315

	@@ -8,8 +8,8 @@ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A
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	-#ifndef TMULA_HPP	11	+#ifndef TMULADDDST_HPP
12	-#define TMULA_HPP	12	+#define TMULADDDST_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
18	namespace pto {	18	namespace pto {
19		19
20	template <typename T>	20	template <typename T>
21	-struct MulaOp {	21	+struct MulAddDstOp {
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 MulaOp {
34	template <	34	template <
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 TMula(	37	+__tf__ PTO_INTERNAL void TMulAddDst(
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 TMula(
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<MulaOp<T>, T, TileDataDst, elementsPerRepeat, blockSizeElem, dstRowStride>(	46	+ BinaryInstr<MulAddDstOp<T>, T, TileDataDst, elementsPerRepeat, blockSizeElem, dstRowStride>(
47	dstPtr, src0Ptr, src1Ptr, validRows, validCols);	47	dstPtr, src0Ptr, src1Ptr, validRows, validCols);
48	} else {	48	} else {
49	- BinaryInstr<MulaOp<T>, T, elementsPerRepeat, blockSizeElem, dstRowStride, src0RowStride, src1RowStride>(	49	+ BinaryInstr<MulAddDstOp<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
55	template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>	55	template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>
56	-PTO_INTERNAL void TMulaCheck(const TileDataDst& dst, const TileDataSrc0& src0, const TileDataSrc1& src1)	56	+PTO_INTERNAL void TMulAddDstCheck(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: TMULA the data type of dst must be consistent with of src0 and src1.");	61	+ "Fix: TMULADDDST 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: TMULA has invalid data type.");	65	+ "Fix: TMULADDDST 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: TMULA only support row major layout.");	68	+ "Fix: TMULADDDST 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: TMULA input tile src0 valid shape mismatch with output tile dst shape.");	74	+ "Fix: TMULADDDST input tile src0 valid shape mismatch with output tile dst shape.");
75	}	75	}
76		76
77	template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>	77	template <typename TileDataDst, typename TileDataSrc0, typename TileDataSrc1>
78	-PTO_INTERNAL void TMULA_IMPL(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1)	78	+PTO_INTERNAL void TMULADDDST_IMPL(TileDataDst& dst, TileDataSrc0& src0, TileDataSrc1& src1)
79	{	79	{
80	using T = typename TileDataDst::DType;	80	using T = typename TileDataDst::DType;
81	- TMulaCheck<TileDataDst, TileDataSrc0, TileDataSrc1>(dst, src0, src1);	81	+ TMulAddDstCheck<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	- TMula<	88	+ TMulAddDst<
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	}

	@@ -15,7 +15,7 @@ using namespace pto;
15	template <	15	template <
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 runTMULA(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1)	18	+__global__ AICORE void runTMULADDDST(__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 runTMULA(__gm__ T __out__* out, __gm__ T __in__* src0, __
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	#endif	49	#endif
50	- TMULA<TileDataDst, TileDataSrc0, TileDataSrc1>(dstTile, src0Tile, src1Tile);	50	+ TMULADDDST<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,22 +59,23 @@ __global__ AICORE void runTMULA(__gm__ T __out__* out, __gm__ T __in__* src0, __
59	template <	59	template <
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 LaunchTMULA(T* out, T* src0, T* src1, void* stream)	62	+void LaunchTMULADDDST(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	- runTMULA<half, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(	65	+ runTMULADDDST<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	- runTMULA<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(out, src0, src1);	68	+ runTMULADDDST<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>(out, src0, src1);
69	}	69	}
70	}	70	}
71		71
72	-template void LaunchTMULA<float, 64, 64, 64, 64, 64, 64, 64, 64>(float* out, float* src0, float* src1, void* stream);	72	+template void LaunchTMULADDDST<float, 64, 64, 64, 64, 64, 64, 64, 64>(
73	-template void LaunchTMULA<float, 32, 128, 32, 192, 32, 256, 32, 127>(
74	float* out, float* src0, float* src1, void* stream);	73	float* out, float* src0, float* src1, void* stream);
75	-template void LaunchTMULA<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(	74	+template void LaunchTMULADDDST<float, 32, 128, 32, 192, 32, 256, 32, 127>(
		75	+ float* out, float* src0, float* src1, void* stream);
		76	+template void LaunchTMULADDDST<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(
76	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);	77	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);
77	-template void LaunchTMULA<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>(	78	+template void LaunchTMULADDDST<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>(
78	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);	79	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);
79	-template void LaunchTMULA<aclFloat16, 1, 16384, 1, 16384, 1, 16384, 1, 16384>(	80	+template void LaunchTMULADDDST<aclFloat16, 1, 16384, 1, 16384, 1, 16384, 1, 16384>(
80	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);	81	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);

	@@ -17,7 +17,7 @@ using namespace pto;
17	template <	17	template <
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 runTMULA(__gm__ T __out__* out, __gm__ T __in__* src0, __gm__ T __in__* src1)	20	+__global__ AICORE void runTMULADDDST(__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 runTMULA(__gm__ T __out__* out, __gm__ T __in__* src0, __
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	#endif	51	#endif
52	- TMULA<TileDataDst, TileDataSrc0, TileDataSrc1>(dstTile, src0Tile, src1Tile);	52	+ TMULADDDST<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,21 +61,22 @@ __global__ AICORE void runTMULA(__gm__ T __out__* out, __gm__ T __in__* src0, __
61	template <	61	template <
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 LaunchTMULA(T* out, T* src0, T* src1, void* stream)	64	+void LaunchTMULADDDST(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	- runTMULA<half, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>	67	+ runTMULADDDST<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	- runTMULA<T, dstTileH, dstTileW, src0TileH, src0TileW, src1TileH, src1TileW, vRows, vCols>	70	+ runTMULADDDST<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 LaunchTMULA<float, 64, 64, 64, 64, 64, 64, 64, 64>(float* out, float* src0, float* src1, void* stream);	75	+template void LaunchTMULADDDST<float, 64, 64, 64, 64, 64, 64, 64, 64>(
76	-template void LaunchTMULA<float, 32, 128, 32, 192, 32, 256, 32, 127>(
77	float* out, float* src0, float* src1, void* stream);	76	float* out, float* src0, float* src1, void* stream);
78	-template void LaunchTMULA<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(	77	+template void LaunchTMULADDDST<float, 32, 128, 32, 192, 32, 256, 32, 127>(
		78	+ float* out, float* src0, float* src1, void* stream);
		79	+template void LaunchTMULADDDST<aclFloat16, 64, 64, 64, 64, 64, 64, 64, 64>(
79	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);	80	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);
80	-template void LaunchTMULA<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>(	81	+template void LaunchTMULADDDST<aclFloat16, 32, 128, 32, 192, 32, 256, 32, 127>(
81	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);	82	aclFloat16* out, aclFloat16* src0, aclFloat16* src1, void* stream);

	@@ -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_32x127	531	python3 tests/script/run_st.py $ARGS -w -v a3 -t tmadd -g TMADDTest.case_float_32x128_32x192_32x256_32x127
532	fi	532	fi
533	if [ -z "$ST_PART" ] \|\| [ "$ST_PART" = "2" ]; then	533	if [ -z "$ST_PART" ] \|\| [ "$ST_PART" = "2" ]; then
534	- python3 tests/script/run_st.py $ARGS -w -v a3 -t tmula -g TMULATest.case_float_32x128_32x192_32x256_32x127	534	+ python3 tests/script/run_st.py $ARGS -w -v a3 -t tmuladddst -g TMULADDDSTTest.case_float_32x128_32x192_32x256_32x127
535	fi	535	fi
536	if [ -z "$ST_PART" ] \|\| [ "$ST_PART" = "2" ]; then	536	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_half	537	python3 tests/script/run_st.py $ARGS -w -v a3 -t textract_ndto2xnz -g TExtractNd2xNzTest.case_half
	@@ -816,7 +816,6 @@ if [ "$ENABLE_A3" = "true" ]; then # A2A3
816	python3 tests/script/run_st.py $ARGS -w -v a3 -t textract_ndto2xnz	816	python3 tests/script/run_st.py $ARGS -w -v a3 -t textract_ndto2xnz
817	python3 tests/script/run_st.py $ARGS -w -v a3 -t tinsert_vec	817	python3 tests/script/run_st.py $ARGS -w -v a3 -t tinsert_vec
818	python3 tests/script/run_st.py $ARGS -w -v a3 -t tmadd	818	python3 tests/script/run_st.py $ARGS -w -v a3 -t tmadd
819	- python3 tests/script/run_st.py $ARGS -w -v a3 -t tmula
820	if [ "$IS_AUTO_MODE" = "false" ]; then	819	if [ "$IS_AUTO_MODE" = "false" ]; then
821	python3 tests/script/run_st.py $ARGS -w -v a3 -t tcolgather	820	python3 tests/script/run_st.py $ARGS -w -v a3 -t tcolgather
822	python3 tests/script/run_st.py $ARGS -w -v a3 -t tcolscatter	821	python3 tests/script/run_st.py $ARGS -w -v a3 -t tcolscatter
	@@ -985,7 +984,6 @@ if [ "$ENABLE_A5" = "true" ]; then
985	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmrgsort -g TMRGSORTTest.case_topk1	984	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmrgsort -g TMRGSORTTest.case_topk1
986	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmul -g TMULTest.case_float_64x64_64x64_64x64_64x64	985	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmul -g TMULTest.case_float_64x64_64x64_64x64_64x64
987	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmadd -g TMADDTest.case_float_32x128_32x192_32x256_32x127	986	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmadd -g TMADDTest.case_float_32x128_32x192_32x256_32x127
988	- python3 tests/script/run_st.py $ARGS -w -v a5 -t tmula -g TMULATest.case_float_32x128_32x192_32x256_32x127
989	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmuls -g TMULSTest.case1	987	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmuls -g TMULSTest.case1
990	python3 tests/script/run_st.py $ARGS -w -v a5 -t tor -g TORTest.case2	988	python3 tests/script/run_st.py $ARGS -w -v a5 -t tor -g TORTest.case2
991	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpow -g TPOWTest.case1	989	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpow -g TPOWTest.case1
	@@ -1162,7 +1160,6 @@ if [ "$ENABLE_A5" = "true" ]; then
1162	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmrgsort	1160	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmrgsort
1163	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmul	1161	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmul
1164	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmadd	1162	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmadd
1165	- python3 tests/script/run_st.py $ARGS -w -v a5 -t tmula
1166	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmuls	1163	python3 tests/script/run_st.py $ARGS -w -v a5 -t tmuls
1167	python3 tests/script/run_st.py $ARGS -w -v a5 -t tor	1164	python3 tests/script/run_st.py $ARGS -w -v a5 -t tor
1168	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpow	1165	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpow