fix docs-asc-devkit-AtomGit

已合并

fix docs #4086

yhkz9211创建于 7月7日从yhkz9211/asc-devkit_9.1.0： sync 合入到 cann/asc-devkit： 9.1.0

fix docs #4086

已合并

yhkz9211创建于 7月7日

从yhkz9211/asc-devkit_9.1.0 sync合入到cann/asc-devkit 9.1.0

收起

classify_rule.yaml
docs / api / SIMD-API
- 基础API
  - Reg矢量计算
    - 数据压缩
      - Unsqueeze.md
    - 直方图计算
      - Histograms.md
  - 矩阵计算（ISASI）
    - Mmad计算
      - Gemm（废弃）.md
      - Mmad.md
      - MmadBitMode.md
      - MmadWithSparse.md
    - 矩阵计算的搬入
      - 矩阵数据搬入至L0-Buffer
        Load2D.md
        Load3D.md
        LoadDataWithTranspose.md
      - 矩阵数据搬入至L1-Buffer
        GMToL1-Load2D指令搬运.md
    - 矩阵计算的搬出
      - L0C到GM数据搬运（DataCopy）.md
      - L0C到GM数据搬运（Fixpipe）.md
      - L0C到L1数据搬运（DataCopy）.md
      - L0C到L1数据搬运（Fixpipe）.md
      - L0C到UB数据搬运（Fixpipe）.md
      - 关键特性说明
        F32-Channel-Split.md
      - 寄存器配置说明
        SetFixPipeAddr.md
        SetFixPipeClipRelu.md
        SetFixPipeConfig.md
        SetFixpipeNz2ndFlag.md
        SetFixpipePreQuantFlag.md
- 高阶API / 张量变换
  - GetBroadCastMaxMinTmpSize.md
  - GetTransDataMaxMinTmpSize.md
  - Pad-Tiling.md
  - Transpose-Tiling.md
  - UnPad-Tiling.md
  - UnPad.md
impl / basic_api
- kernel_operator_gemm_base_impl.h
- kernel_operator_gemm_intf_impl.h
include / basic_api
- kernel_operator_gemm_intf.h
tests / api / basic_api
- ascendc_case_common
  - test_operator_gemm.cpp
- ascendc_header_checker
  - kernel_operator_gemm_intf.cpp

共 33 个文件变更+283-328

Mclassify_rule.yaml+2-0

	@@ -60,6 +60,8 @@ asc-devkit:
60	- impl/c_api/instr_impl/npu_arch_2201/debug_impl/printf_impl.h	60	- impl/c_api/instr_impl/npu_arch_2201/debug_impl/printf_impl.h
61	- impl/c_api/instr_impl/npu_arch_2201/debug_impl.h	61	- impl/c_api/instr_impl/npu_arch_2201/debug_impl.h
62	- impl/basic_api/common_types.h	62	- impl/basic_api/common_types.h
		63	+ - impl/basic_api/kernel_operator_gemm_base_impl.h
		64	+ - impl/basic_api/kernel_operator_gemm_intf_impl.h
63	- impl/adv_api/detail/hccl/common/hccl_inner_def.h	65	- impl/adv_api/detail/hccl/common/hccl_inner_def.h
64	- impl/adv_api/detail/hccl/common/hccl_base.h	66	- impl/adv_api/detail/hccl/common/hccl_base.h
65	- impl/adv_api/detail/hccl/common/hccl_aicpu_def.h	67	- impl/adv_api/detail/hccl/common/hccl_aicpu_def.h

Mdocs/api/SIMD-API/基础API/Reg矢量计算/数据压缩/Unsqueeze.md+14-30

	@@ -49,12 +49,12 @@
49		49
50	## 功能说明<a name="section618mcpsimp"></a>	50	## 功能说明<a name="section618mcpsimp"></a>
51		51
52	-将dstReg中数据根据mask进行解压缩。解压缩方式：dstReg中第0个元素置为0，dstReg中的第i个元素等于mask中从第0个到第\(i-1\)个元素中1的数量。mask最高位被忽略不参与统计。以uint8\_t类型为例，算法逻辑表示如下：	52	+以dstReg为源操作数和目的操作数，根据mask进行解压缩。解压缩方式：dstReg中第0个元素置为0，dstReg中的第i个元素等于mask中从第0个到第\(i-1\)个元素中1的数量。mask最高位被忽略不参与统计。以uint8\_t类型为例，算法逻辑表示如下：
53		53
54	```	54	```
55	dstReg[0] = 0;	55	dstReg[0] = 0;
56	for(int i = 1; i < AscendC::GetVecLen() / sizeof(uint8_t); i++){	56	for(int i = 1; i < AscendC::GetVecLen() / sizeof(uint8_t); i++){
57	-dstReg[i] = mask[i - 1] ? (dstReg[i - 1] + 1) : dstReg[i - 1];	57	+ dstReg[i] = mask[i - 1] ? (dstReg[i - 1] + 1) : dstReg[i - 1];
58	}	58	}
59	```	59	```
60		60
	@@ -94,33 +94,18 @@ __simd_callee__ inline void Unsqueeze(U& dstReg, MaskReg& mask)
94		94
95	表2 函数参数说明	95	表2 函数参数说明
96		96
97	-<a name="zh-cn_topic_0235751031_table33761356"></a>	97	+\| 参数名称 \| 输入/输出 \| 描述 \|
		Y Y Yifan Wang7月7日 [V0-P2] 文档清理 \| Unsqueeze参数描述需同步 dstReg参数标注已从"输出"改为"输入/输出"（正确，因为Unsqueeze以dstReg为源和目的），但Histograms.md中dstReg参数表未同步清理为markdown格式（仍为HTML table）。建议确认是否遗漏。此评论由AI生成，如有异议，请回复帮助我更快成长：）
98	-<table><thead align="left"><tr id="zh-cn_topic_0235751031_row27598891"><th class="cellrowborder" valign="top" width="18.54%" id="mcps1.2.4.1.1"><p id="zh-cn_topic_0235751031_p20917673"><a name="zh-cn_topic_0235751031_p20917673"></a><a name="zh-cn_topic_0235751031_p20917673"></a>参数名</p>	98	+\| ------ \| ------ \| ------ \|
99	-</th>	99	+\| dstReg \| 输入/输出 \| 源操作数和目的操作数。<br>类型为[RegTensor](../寄存器数据类型/RegTensor.md)。 \|
100	-<th class="cellrowborder" valign="top" width="10.05%" id="mcps1.2.4.1.2"><p id="zh-cn_topic_0235751031_p16609919"><a name="zh-cn_topic_0235751031_p16609919"></a><a name="zh-cn_topic_0235751031_p16609919"></a>输入/输出</p>	100	+\| mask \| 输入 \| mask用于提供dstReg解压缩信息。<br>类型为[MaskReg](../寄存器数据类型/MaskReg.md)。 \|
101	-</th>	101	+
102	-<th class="cellrowborder" valign="top" width="71.41%" id="mcps1.2.4.1.3"><p id="zh-cn_topic_0235751031_p59995477"><a name="zh-cn_topic_0235751031_p59995477"></a><a name="zh-cn_topic_0235751031_p59995477"></a>描述</p>	102	+## 返回值说明
103	-</th>	103	+
104	-</tr>	104	+无
105	-</thead>	105	+
106	-<tbody><tr id="row42461942101815"><td class="cellrowborder" valign="top" width="18.54%" headers="mcps1.2.4.1.1 "><p id="p284425844311"><a name="p284425844311"></a><a name="p284425844311"></a>dstReg</p>	106	+## 数据类型
107	-</td>	107	+
108	-<td class="cellrowborder" valign="top" width="10.05%" headers="mcps1.2.4.1.2 "><p id="p158449584436"><a name="p158449584436"></a><a name="p158449584436"></a>输出</p>	108	+支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t。
109	-</td>
110	-<td class="cellrowborder" valign="top" width="71.41%" headers="mcps1.2.4.1.3 "><p id="p6389114616514"><a name="p6389114616514"></a><a name="p6389114616514"></a>目的操作数和源操作数。</p>
111	-<p id="p1818424418593"><a name="p1818424418593"></a><a name="p1818424418593"></a>类型为<a href="../寄存器数据类型/RegTensor.md">RegTensor</a>。</p>
112	-</td>
113	-</tr>
114	-<tr id="row17332152321719"><td class="cellrowborder" valign="top" width="18.54%" headers="mcps1.2.4.1.1 "><p id="p1538745616196"><a name="p1538745616196"></a><a name="p1538745616196"></a>mask</p>
115	-</td>
116	-<td class="cellrowborder" valign="top" width="10.05%" headers="mcps1.2.4.1.2 "><p id="p73331523101710"><a name="p73331523101710"></a><a name="p73331523101710"></a>输入</p>
117	-</td>
118	-<td class="cellrowborder" valign="top" width="71.41%" headers="mcps1.2.4.1.3 "><p id="p1559991025517"><a name="p1559991025517"></a><a name="p1559991025517"></a>mask用于提供dstReg解压缩信息。</p>
119	-<p id="p911252842816"><a name="p911252842816"></a><a name="p911252842816"></a>类型为MaskReg。</p>
120	-</td>
121	-</tr>
122	-</tbody>
123	-</table>
124		109
125	## 约束说明<a name="section177921451558"></a>	110	## 约束说明<a name="section177921451558"></a>
126		111
	@@ -141,4 +126,3 @@ __simd_vf__ inline void UnsqueezeVF(__ubuf__ T* dstAddr, uint32_t oneRepeatSize,
141	}	126	}
142	}	127	}
143	```	128	```
144	-

Mdocs/api/SIMD-API/基础API/Reg矢量计算/直方图计算/Histograms.md+8-43

	@@ -76,49 +76,14 @@ __simd_callee__ inline void Histograms(V& dstReg, S& srcReg, MaskReg& mask)
76		76
77	表1 模板参数说明	77	表1 模板参数说明
78		78
79	-<a name="table4835205712588"></a>	79	+\| 参数名称 \| 描述 \|
80	-<table><thead align="left"><tr id="row118356578583"><th class="cellrowborder" valign="top" width="17.94%" id="mcps1.2.3.1.1"><p id="p48354572582"><a name="p48354572582"></a><a name="p48354572582"></a>参数名</p>	80	+\| ------ \| ------ \|
81	-</th>	81	+\| T \| 源操作数的数据类型。 \|
82	-<th class="cellrowborder" valign="top" width="82.06%" id="mcps1.2.3.1.2"><p id="p583535795817"><a name="p583535795817"></a><a name="p583535795817"></a>描述</p>	82	+\| U \| 目的操作数的数据类型。 \|
83	-</th>	83	+\| mode \| HistogramsBinType枚举类型，用于选择统计srcReg的低位区间或高位区间。Histograms将uint8_t类型srcReg的[0, 255]取值范围划分为低位、高位两个区间；每个区间包含128个取值，对应uint16_t类型dstReg在一个VL（256Byte）内可保存的128个统计结果。单次调用只统计其中一个区间，并将该区间的统计结果写入传入的dstReg。若需要获得完整[0, 255]范围的统计结果，需要分别以BIN0和BIN1调用两次接口，并使用两个目的寄存器保存结果，例如dst0保存低位区间结果，dst1保存高位区间结果。<br>• BIN0：统计srcReg中[0, 127]范围内的数据，结果写入传入的dstReg。<br>• BIN1：统计srcReg中[128, 255]范围内的数据，结果写入传入的dstReg。 \|
84	-</tr>	84	+\| type \| HistogramsType 枚举类型，表示统计模式。<br>• FREQUENCY：频率统计模式，统计srcReg中[0, 255]每个数的数量。每个dst有128个元素，其中低位模式dstReg中每个元素对应src中[0, 127]每个元素的累加个数，高位模式dstReg中每个元素对应src中[128,255]每个元素的累加个数。<br>• ACCUMULATE：累计统计模式，统计srcReg中[0, 255]每个数及其之前出现数的数量总和。每个dst有128个元素，其中低位模式dstReg中每个元素对应src中[0, 127]每个元素区间累加个数，高位模式dstReg中每个元素对应src中[128,255]每个元素区间累加个数。 \|
85	-</thead>	85	+\| S \| 源操作数RegTensor类型，由编译器自动推导，用户不需要填写。 \|
86	-<tbody><tr id="row1835857145817"><td class="cellrowborder" valign="top" width="17.94%" headers="mcps1.2.3.1.1 "><p id="p5835457165816"><a name="p5835457165816"></a><a name="p5835457165816"></a>T</p>	86	+\| V \| 目的操作数RegTensor类型，由编译器自动推导，用户不需要填写。 \|
87	-</td>
88	-<td class="cellrowborder" valign="top" width="82.06%" headers="mcps1.2.3.1.2 "><p id="p168351657155818"><a name="p168351657155818"></a><a name="p168351657155818"></a>源操作数的数据类型。</p>
89	-<p id="p188929596549"><a name="p188929596549"></a><a name="p188929596549"></a><span id="ph1479120155519"><a name="ph1479120155519"></a><a name="ph1479120155519"></a>Ascend 950PR/Ascend 950DT</span>，支持的数据类型为：uint8_t</p>
90	-</td>
91	-</tr>
92	-<tr id="row18835145716587"><td class="cellrowborder" valign="top" width="17.94%" headers="mcps1.2.3.1.1 "><p id="p1383515717581"><a name="p1383515717581"></a><a name="p1383515717581"></a>U</p>
93	-</td>
94	-<td class="cellrowborder" valign="top" width="82.06%" headers="mcps1.2.3.1.2 "><p id="p1775118537244"><a name="p1775118537244"></a><a name="p1775118537244"></a>目的操作数的数据类型。</p>
95	-<p id="p18114175555419"><a name="p18114175555419"></a><a name="p18114175555419"></a><span id="ph811495510548"><a name="ph811495510548"></a><a name="ph811495510548"></a>Ascend 950PR/Ascend 950DT</span>，支持的数据类型为：uint16_t</p>
96	-</td>
97	-</tr>
98	-<tr id="row18128134152619"><td class="cellrowborder" valign="top" width="17.94%" headers="mcps1.2.3.1.1 "><p id="p2012917416262"><a name="p2012917416262"></a><a name="p2012917416262"></a>mode</p>
99	-</td>
100	-<td class="cellrowborder" valign="top" width="82.06%" headers="mcps1.2.3.1.2 "><p id="p9905165812547"><a name="p9905165812547"></a><a name="p9905165812547"></a>HistogramsBinType枚举类型，用于控制统计src的低位区间[0-127]（前半部分）或高位区间[128-255]（后半部分）的数据。VL长度为256Byte，dst数据类型为uint16_t，一个dst可以存储128个数据，因此需要两个dst。</p>
101	-<a name="ul59931348346"></a><a name="ul59931348346"></a><ul id="ul59931348346"><li>BIN0表示低位模式，统计src中[0-127]范围内的数据写入dst0。</li><li>BIN1表示高位模式，统计src中[128-255]范围内的数据写入dst1。</li></ul>
102	-</td>
103	-</tr>
104	-<tr id="row134244441918"><td class="cellrowborder" valign="top" width="17.94%" headers="mcps1.2.3.1.1 "><p id="p916243141912"><a name="p916243141912"></a><a name="p916243141912"></a>type</p>
105	-</td>
106	-<td class="cellrowborder" valign="top" width="82.06%" headers="mcps1.2.3.1.2 "><div class="p" id="p1663311238518"><a name="p1663311238518"></a><a name="p1663311238518"></a>HistogramsType枚举类型，表示统计模式。<a name="ul36112038046"></a><a name="ul36112038046"></a><ul id="ul36112038046"><li>FREQUENCY：频率统计模式，统计src中[0, 255]每个数的数量。每个dst有128个元素，其中dst0中每个元素对应src中[0, 127]每个元素的累加个数，dst1中每个元素对应src中[128,255]每个元素的累加个数。</li><li>ACCUMULATE：累计统计模式，统计src中x<=0、x<=1、x<=2、x<=3......x<=254、x<=255每个区间的元素个数。每个dst有128个元素，其中dst0中每个元素对应src中[0, 127]每个元素区间累加个数，dst1中每个元素对应src中[128，255]每个元素区间累加个数。</li></ul>
107	-</div>
108	-</td>
109	-</tr>
110	-<tr id="row1237743311380"><td class="cellrowborder" valign="top" width="17.94%" headers="mcps1.2.3.1.1 "><p id="p6377173353817"><a name="p6377173353817"></a><a name="p6377173353817"></a>S</p>
111	-</td>
112	-<td class="cellrowborder" valign="top" width="82.06%" headers="mcps1.2.3.1.2 "><p id="p2304311103410"><a name="p2304311103410"></a><a name="p2304311103410"></a>目的操作数RegTensor类型，由编译器自动推导，用户不需要填写。</p>
113	-</td>
114	-</tr>
115	-<tr id="row175751956133812"><td class="cellrowborder" valign="top" width="17.94%" headers="mcps1.2.3.1.1 "><p id="p357565618387"><a name="p357565618387"></a><a name="p357565618387"></a>V</p>
116	-</td>
117	-<td class="cellrowborder" valign="top" width="82.06%" headers="mcps1.2.3.1.2 "><p id="p12801441174416"><a name="p12801441174416"></a><a name="p12801441174416"></a>源操作数RegTensor类型，由编译器自动推导，用户不需要填写。</p>
118	-</td>
119	-</tr>
120	-</tbody>
121	-</table>
122		87
123	表2 参数说明	88	表2 参数说明
124		89

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/Mmad计算/Gemm（废弃）.md+4-4

	@@ -36,7 +36,7 @@
36		36
37	```cpp	37	```cpp
38	template <typename T, typename U, typename S>	38	template <typename T, typename U, typename S>
39	- __aicore__ inline void Gemm(const LocalTensor<T>& dst, const LocalTensor<U>& src0, const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling tilling, bool partialsum = true, int32_t initValue = 0)	39	+ __aicore__ inline void Gemm(const LocalTensor<T>& dst, const LocalTensor<U>& src0, const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling tiling, bool partialsum = true, int32_t initValue = 0)
40	```	40	```
41		41
42	- 切分方案计算接口：	42	- 切分方案计算接口：
	@@ -58,9 +58,9 @@
58	\| m \| 输入 \| 左矩阵Src0Local有效Height，范围：[1, 4096]。<br>注意：m可以不是16的倍数。 \|	58	\| m \| 输入 \| 左矩阵Src0Local有效Height，范围：[1, 4096]。<br>注意：m可以不是16的倍数。 \|
59	\| k \| 输入 \| 左矩阵Src0Local有效Width、右矩阵Src1Local有效Height。<br>• 当输入张量Src0Local的数据类型为float时，范围：[1, 8192]<br>• 当输入张量Src0Local的数据类型为half时，范围：[1, 16384]<br>• 当输入张量Src0Local的数据类型为int8_t时，范围：[1, 32768]<br><br>注意：k可以不是16的倍数。 \|	59	\| k \| 输入 \| 左矩阵Src0Local有效Width、右矩阵Src1Local有效Height。<br>• 当输入张量Src0Local的数据类型为float时，范围：[1, 8192]<br>• 当输入张量Src0Local的数据类型为half时，范围：[1, 16384]<br>• 当输入张量Src0Local的数据类型为int8_t时，范围：[1, 32768]<br><br>注意：k可以不是16的倍数。 \|
60	\| n \| 输入 \| 右矩阵Src1Local有效Width，范围：[1, 4096]。<br>注意：n可以不是16的倍数。 \|	60	\| n \| 输入 \| 右矩阵Src1Local有效Width，范围：[1, 4096]。<br>注意：n可以不是16的倍数。 \|
61	-\| tilling \| 输入 \| 切分规则，类型为GemmTiling，结构体具体定义为：<br><br><pre>struct GemmTiling {<br> const uint32_t blockSize = 16;<br> LoopMode loopMode = LoopMode::MODE_NM;<br> uint32_t mNum = 0;<br> uint32_t nNum = 0;<br> uint32_t kNum = 0;<br> uint32_t roundM = 0;<br> uint32_t roundN = 0;<br> uint32_t roundK = 0;<br> uint32_t c0Size = 32;<br> uint32_t dtypeSize = 1;<br> uint32_t mBlockNum = 0;<br> uint32_t nBlockNum = 0;<br> uint32_t kBlockNum = 0;<br> uint32_t mIterNum = 0;<br> uint32_t nIterNum = 0;<br> uint32_t kIterNum = 0;<br> uint32_t mTileBlock = 0;<br> uint32_t nTileBlock = 0;<br> uint32_t kTileBlock = 0;<br> uint32_t kTailBlock = 0;<br> uint32_t mTailBlock = 0;<br> uint32_t nTailBlock = 0;<br> bool kHasTail = false;<br> bool mHasTail = false;<br> bool nHasTail = false;<br> bool kHasTailEle = false;<br> uint32_t kTailEle = 0;<br>};<br></pre><br>参数说明请参考[表3](#table3)。 \|	61	+\| tiling \| 输入 \| 切分规则，类型为GemmTiling，结构体具体定义为：<br><br><pre>struct GemmTiling {<br> const uint32_t blockSize = 16;<br> LoopMode loopMode = LoopMode::MODE_NM;<br> uint32_t mNum = 0;<br> uint32_t nNum = 0;<br> uint32_t kNum = 0;<br> uint32_t roundM = 0;<br> uint32_t roundN = 0;<br> uint32_t roundK = 0;<br> uint32_t c0Size = 32;<br> uint32_t dtypeSize = 1;<br> uint32_t mBlockNum = 0;<br> uint32_t nBlockNum = 0;<br> uint32_t kBlockNum = 0;<br> uint32_t mIterNum = 0;<br> uint32_t nIterNum = 0;<br> uint32_t kIterNum = 0;<br> uint32_t mTileBlock = 0;<br> uint32_t nTileBlock = 0;<br> uint32_t kTileBlock = 0;<br> uint32_t kTailBlock = 0;<br> uint32_t mTailBlock = 0;<br> uint32_t nTailBlock = 0;<br> bool kHasTail = false;<br> bool mHasTail = false;<br> bool nHasTail = false;<br> bool kHasTailEle = false;<br> uint32_t kTailEle = 0;<br>};<br></pre><br>参数说明请参考[表3](#table3)。 \|
62	-\| partialsum \| 输入 \| 当dst参数所在的TPosition为CO2时，通过该参数控制计算结果是否搬出。<br>• 取值0：搬出计算结果；<br>• 取值1：不搬出计算结果，可以进行后续计算。 \|	62	+\| partialsum \| 输入 \| 当dst参数所在的TPosition为CO2时，通过该参数控制计算结果是否搬出。<br>•取值0：搬出计算结果；<br>•取值1：不搬出计算结果，可以进行后续计算。 \|
63	-\| initValue \| 输入 \| 表示dst是否需要初始化。<br>• 取值0：dst需要初始化，dst初始矩阵保存有之前结果，新计算结果会累加前一次Gemm计算结果。<br>• 取值1：dst不需要初始化，dst初始矩阵中数据无意义，计算结果直接覆盖dst中的数据。 \|	63	+\| initValue \| 输入 \| 表示dst是否需要初始化。<br>•取值0：dst需要初始化，dst初始矩阵保存有之前结果，新计算结果会累加前一次Gemm计算结果。<br>•取值1：dst不需要初始化，dst初始矩阵中数据无意义，计算结果直接覆盖dst中的数据。 \|
64		64
65	表2 src0、src1和dst的数据类型组合	65	表2 src0、src1和dst的数据类型组合
66		66

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/Mmad计算/Mmad.md+4-4

	@@ -164,10 +164,10 @@ Atlas A2 训练系列产品/Atlas A2 推理系列产品
164		164
165	\| 参数名称 \| 含义 \|	165	\| 参数名称 \| 含义 \|
166	\| --- \| --- \|	166	\| --- \| --- \|
167	-\| m \| 左矩阵Height，取值范围：m∈[0，4095]。默认值为0。 \|	167	+\| m \| 左矩阵Height，取值范围：m∈[0, 4095]。默认值为0。 \|
168	-\| n \| 右矩阵Width，取值范围：n∈[0，4095]。默认值为0。 \|	168	+\| n \| 右矩阵Width，取值范围：n∈[0, 4095]。默认值为0。 \|
169	-\| k \| 左矩阵Width、右矩阵Height，取值范围：k∈[0，4095]。默认值为0。 \|	169	+\| k \| 左矩阵Width、右矩阵Height，取值范围：k∈[0, 4095]。默认值为0。 \|
170	-\| cmatrixInitVal \| 是否开启C矩阵默认初始化清零操作。默认值true。<br>  • true：C矩阵默认初始化为0；<br>  • false：C矩阵不进行默认操作，通过设置cmatrixSource参数进行初始化。 \|	170	+\| cmatrixInitVal \| 是否开启C矩阵默认初始化清零操作。默认值为 true。<br>  • true：C矩阵默认初始化为0；<br>  • false：C矩阵不进行默认操作，通过设置cmatrixSource参数进行初始化。 \|
171	\| cmatrixSource \| 配置C矩阵初始值是否来源于BT Buffer。默认值为false。<br>  • false：不对L0C Buffer进行初始化操作；<br>  • true：使用BT Buffer(TPosition:C2)的数据对L0C Buffer进行初始化操作。<br><br>Atlas 训练系列产品，仅支持配置为false。<br><br>Atlas 推理系列产品AI Core，仅支持配置为false。<br><br>Atlas A2 训练系列产品/Atlas A2 推理系列产品，支持配置为true/false。<br><br>Atlas A3 训练系列产品/Atlas A3 推理系列产品，支持配置为true/false。<br><br>Atlas 200I/500 A2 推理产品，支持配置为true/false。<br><br>Ascend 950PR/Ascend 950DT，支持配置为true/false。<br><br>Kirin X90仅支持配置为false。<br><br>Kirin 9030仅支持配置为false。<br><br>注意：带Bias输入的接口配置该参数无效，会根据bias输入的位置来判断C矩阵初始值是否来源于BT Buffer。 \|	171	\| cmatrixSource \| 配置C矩阵初始值是否来源于BT Buffer。默认值为false。<br>  • false：不对L0C Buffer进行初始化操作；<br>  • true：使用BT Buffer(TPosition:C2)的数据对L0C Buffer进行初始化操作。<br><br>Atlas 训练系列产品，仅支持配置为false。<br><br>Atlas 推理系列产品AI Core，仅支持配置为false。<br><br>Atlas A2 训练系列产品/Atlas A2 推理系列产品，支持配置为true/false。<br><br>Atlas A3 训练系列产品/Atlas A3 推理系列产品，支持配置为true/false。<br><br>Atlas 200I/500 A2 推理产品，支持配置为true/false。<br><br>Ascend 950PR/Ascend 950DT，支持配置为true/false。<br><br>Kirin X90仅支持配置为false。<br><br>Kirin 9030仅支持配置为false。<br><br>注意：带Bias输入的接口配置该参数无效，会根据bias输入的位置来判断C矩阵初始值是否来源于BT Buffer。 \|
172	\| isBias \| 该参数废弃，新开发内容不要使用该参数。如果需要累加初始矩阵，请使用带Bias的接口来实现；也可以通过cmatrixInitVal和cmatrixSource参数配置C矩阵的初始值来源来实现。推荐使用带Bias的接口，相比于配置cmatrixInitVal和cmatrixSource参数更加简单方便。<br><br>配置是否需要累加初始矩阵，默认值为false，取值说明如下：<br>  • false：矩阵乘，无需累加初始矩阵，C = A \* B。<br>  • true：矩阵乘加，需要累加初始矩阵，C += A \* B。 \|	172	\| isBias \| 该参数废弃，新开发内容不要使用该参数。如果需要累加初始矩阵，请使用带Bias的接口来实现；也可以通过cmatrixInitVal和cmatrixSource参数配置C矩阵的初始值来源来实现。推荐使用带Bias的接口，相比于配置cmatrixInitVal和cmatrixSource参数更加简单方便。<br><br>配置是否需要累加初始矩阵，默认值为false，取值说明如下：<br>  • false：矩阵乘，无需累加初始矩阵，C = A \* B。<br>  • true：矩阵乘加，需要累加初始矩阵，C += A \* B。 \|
173	\| disableGemv \| M=1时，该参数用来配置Mmad计算是否开启[GEMV](关键特性说明/GEMV.md#ZH-CN_TOPIC_0000002538231187)模式。<br>  • false：开启GEMV模式。<br>  • true：关闭GEMV模式。<br><br>该参数仅支持如下型号：<br><br>Ascend 950PR/Ascend 950DT \|	173	\| disableGemv \| M=1时，该参数用来配置Mmad计算是否开启[GEMV](关键特性说明/GEMV.md#ZH-CN_TOPIC_0000002538231187)模式。<br>  • false：开启GEMV模式。<br>  • true：关闭GEMV模式。<br><br>该参数仅支持如下型号：<br><br>Ascend 950PR/Ascend 950DT \|

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/Mmad计算/MmadBitMode.md+4-4

	@@ -127,13 +127,13 @@ __aicore__ inline MmadBitModeParams(const MmadBitModeParams &mmadParams_);
127		127
128	\| 参数名称 \| 含义 \|	128	\| 参数名称 \| 含义 \|
129	\| -------------- \| ------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- \|	129	\| -------------- \| ------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- \|
130	-\| m \| 左矩阵Height，取值范围：m∈[0，4095]。默认值为0。<br>该参数是位域结构体的最低位参数，占用12bit，可以使用MmadBitModeParams类对象的SetM()函数设置其值，使用GetM()函数获取其值。 \|	130	+\| m \| 左矩阵Height，取值范围：m∈[0, 4095]。默认值为0。<br>该参数是位域结构体的最低位参数，占用12bit，可以使用MmadBitModeParams类对象的SetM()函数设置其值，使用GetM()函数获取其值。 \|
131	-\| k \| 左矩阵Width、右矩阵Height，取值范围：k∈[0，4095]。默认值为0。<br>该参数是位域结构体的第二低位参数，占用12bit，可以使用MmadBitModeParams类对象的SetK()函数设置其值，使用GetK()函数获取其值。 \|	131	+\| k \| 左矩阵Width、右矩阵Height，取值范围：k∈[0, 4095]。默认值为0。<br>该参数是位域结构体的第二低位参数，占用12bit，可以使用MmadBitModeParams类对象的SetK()函数设置其值，使用GetK()函数获取其值。 \|
132	-\| n \| 右矩阵Width，取值范围：n∈[0，4095]。默认值为0。<br>该参数是位域结构体的第三低位参数，占用12bit，可以使用MmadBitModeParams类对象的SetN()函数设置其值，使用GetN()函数获取其值。 \|	132	+\| n \| 右矩阵Width，取值范围：n∈[0, 4095]。默认值为0。<br>该参数是位域结构体的第三低位参数，占用12bit，可以使用MmadBitModeParams类对象的SetN()函数设置其值，使用GetN()函数获取其值。 \|
133	\| unitFlag \| 预留参数。为后续的功能做保留，开发者暂时无需关注，使用默认值即可。<br>该参数是位域结构体的第四低位参数，占用2bit，可以使用MmadBitModeParams类对象的SetUnitFlag()函数设置其值，使用GetUnitFlag()函数获取其值。 \|	133	\| unitFlag \| 预留参数。为后续的功能做保留，开发者暂时无需关注，使用默认值即可。<br>该参数是位域结构体的第四低位参数，占用2bit，可以使用MmadBitModeParams类对象的SetUnitFlag()函数设置其值，使用GetUnitFlag()函数获取其值。 \|
134	\| disableGemv \| M = 1时，用于配置Mmad计算是否开启GEMV。当输入为false时，表示开启GEMV；反之，输入为true时，表示关闭GEMV。<br>GEMV（General Matrix-Vector Multiplication）表示实现矩阵和向量的乘积，开启GEMV后，Mmad API从L0A Buffer读取数据时，数据将以ND格式进行读取，而不会将其视为ZZ格式。<br>该参数是位域结构体的第五低位参数，占用1bit，可以使用MmadBitModeParams类对象的SetDisableGemv()函数设置其值，使用GetDisableGemv()函数获取其值。 \|	134	\| disableGemv \| M = 1时，用于配置Mmad计算是否开启GEMV。当输入为false时，表示开启GEMV；反之，输入为true时，表示关闭GEMV。<br>GEMV（General Matrix-Vector Multiplication）表示实现矩阵和向量的乘积，开启GEMV后，Mmad API从L0A Buffer读取数据时，数据将以ND格式进行读取，而不会将其视为ZZ格式。<br>该参数是位域结构体的第五低位参数，占用1bit，可以使用MmadBitModeParams类对象的SetDisableGemv()函数设置其值，使用GetDisableGemv()函数获取其值。 \|
135	\| cmatrixSource \| 配置C矩阵初始值是否来源于BT Buffer（TPosition:C2）。默认值为false。<br>  • false：来源于L0C Buffer（TPosition:CO1）；<br>  • true：来源于BT Buffer（TPosition:C2）。<br>注意：带bias输入的接口配置该参数无效，会根据bias输入的位置来判断C矩阵初始值是否来源于L0C Buffer还是BT Buffer。<br>该参数是位域结构体的第六低位参数，占用1bit，可以使用MmadBitModeParams类对象的SetCmatrixSource()函数设置其值，使用GetCmatrixSource()函数获取其值。 \|	135	\| cmatrixSource \| 配置C矩阵初始值是否来源于BT Buffer（TPosition:C2）。默认值为false。<br>  • false：来源于L0C Buffer（TPosition:CO1）；<br>  • true：来源于BT Buffer（TPosition:C2）。<br>注意：带bias输入的接口配置该参数无效，会根据bias输入的位置来判断C矩阵初始值是否来源于L0C Buffer还是BT Buffer。<br>该参数是位域结构体的第六低位参数，占用1bit，可以使用MmadBitModeParams类对象的SetCmatrixSource()函数设置其值，使用GetCmatrixSource()函数获取其值。 \|
136	-\| cmatrixInitVal \| 配置C矩阵初始值是否为0。默认值true。<br>  • true：C矩阵初始值为0；<br>  • false：C矩阵初始值通过cmatrixSource参数进行配置。<br>该参数是位域结构体的最高位参数，占用1bit，可以使用MmadBitModeParams类对象的SetCmatrixInitVal()函数设置其值，使用GetCmatrixInitVal()函数获取其值。 \|	136	+\| cmatrixInitVal \| 配置C矩阵初始值是否为0。默认值为 true。<br>  • true：C矩阵初始值为0；<br>  • false：C矩阵初始值通过cmatrixSource参数进行配置。<br>该参数是位域结构体的最高位参数，占用1bit，可以使用MmadBitModeParams类对象的SetCmatrixInitVal()函数设置其值，使用GetCmatrixInitVal()函数获取其值。 \|
137		137
138	## 数据类型<a id="zh_cn_topic_mmadbitmode_section_datatype"></a>	138	## 数据类型<a id="zh_cn_topic_mmadbitmode_section_datatype"></a>
139		139

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/Mmad计算/MmadWithSparse.md+4-4

	@@ -113,10 +113,10 @@ __aicore__ inline void MmadWithSparse(const LocalTensor<T>& dst, const LocalTens
113		113
114	\| 参数名称 \| 含义 \|	114	\| 参数名称 \| 含义 \|
115	\| --------------- \| --------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- \|	115	\| --------------- \| --------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------- \|
116	-\| m \| 左矩阵Height，取值范围：m∈[0，4095]。默认值为0。 \|	116	+\| m \| 左矩阵Height，取值范围：m∈[0, 4095]。默认值为0。 \|
117	-\| n \| 右矩阵Width，取值范围：n∈[0，4095]。默认值为0。 \|	117	+\| n \| 右矩阵Width，取值范围：n∈[0, 4095]。默认值为0。 \|
118	-\| k \| 左矩阵Width、右矩阵Height，取值范围：k∈[0，4095]。默认值为0。 \|	118	+\| k \| 左矩阵Width、右矩阵Height，取值范围：k∈[0, 4095]。默认值为0。 \|
119	-\| cmatrixInitVal \| 是否开启C矩阵默认初始化清零操作。默认值true。<br>  • true：C矩阵默认初始化为0；<br>  • false：C矩阵不进行默认操作，通过设置cmatrixSource参数进行初始化。 \|	119	+\| cmatrixInitVal \| 是否开启C矩阵默认初始化清零操作。默认值为 true。<br>  • true：C矩阵默认初始化为0；<br>  • false：C矩阵不进行默认操作，通过设置cmatrixSource参数进行初始化。 \|
120	\| cmatrixSource \| 配置C矩阵初始值是否来源于BT Buffer。默认值为false。<br>  • false：不对L0C Buffer进行初始化操作；<br>  • true：使用BT Buffer（TPosition:C2）的数据对L0C Buffer进行初始化操作。<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品，支持配置为true/false。<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品，支持配置为true/false。<br>Atlas 200I/500 A2 推理产品，支持配置为true/false。<br>注意：带Bias输入的接口配置该参数无效，会根据Bias输入的位置来判断C矩阵初始值是否来源于BT Buffer。 \|	120	\| cmatrixSource \| 配置C矩阵初始值是否来源于BT Buffer。默认值为false。<br>  • false：不对L0C Buffer进行初始化操作；<br>  • true：使用BT Buffer（TPosition:C2）的数据对L0C Buffer进行初始化操作。<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品，支持配置为true/false。<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品，支持配置为true/false。<br>Atlas 200I/500 A2 推理产品，支持配置为true/false。<br>注意：带Bias输入的接口配置该参数无效，会根据Bias输入的位置来判断C矩阵初始值是否来源于BT Buffer。 \|
121	\| isBias \| 该参数废弃，新开发内容不要使用该参数。如果需要累加初始矩阵，请使用带Bias的接口来实现；也可以通过cmatrixInitVal和cmatrixSource参数配置C矩阵的初始值来源来实现。推荐使用带Bias的接口，相比于配置cmatrixInitVal和cmatrixSource参数更加简单方便。<br>配置是否需要累加初始矩阵，默认值为false，取值说明如下：<br>  • false：矩阵乘，无需累加初始矩阵，C = A \* B。<br>  • true：矩阵乘加，需要累加初始矩阵，C += A \* B。 \|	121	\| isBias \| 该参数废弃，新开发内容不要使用该参数。如果需要累加初始矩阵，请使用带Bias的接口来实现；也可以通过cmatrixInitVal和cmatrixSource参数配置C矩阵的初始值来源来实现。推荐使用带Bias的接口，相比于配置cmatrixInitVal和cmatrixSource参数更加简单方便。<br>配置是否需要累加初始矩阵，默认值为false，取值说明如下：<br>  • false：矩阵乘，无需累加初始矩阵，C = A \* B。<br>  • true：矩阵乘加，需要累加初始矩阵，C += A \* B。 \|
122	\| unitFlag \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag；<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位；<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>该参数仅支持如下型号：<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品；<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品。<br>参数设置方案和特性细节可参考：[UnitFlag](关键特性说明/UnitFlag.md#ZH-CN_TOPIC_00000025690709788)。 \|	122	\| unitFlag \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag；<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位；<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>该参数仅支持如下型号：<br>Atlas A2 训练系列产品/Atlas A2 推理系列产品；<br>Atlas A3 训练系列产品/Atlas A3 推理系列产品。<br>参数设置方案和特性细节可参考：[UnitFlag](关键特性说明/UnitFlag.md#ZH-CN_TOPIC_00000025690709788)。 \|

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬入/矩阵数据搬入至L0-Buffer/Load2D.md+1-1

	@@ -88,7 +88,7 @@ __aicore__ inline void LoadData(const LocalTensor<T>& dst, const LocalTensor<T>&
88	\| sid \| 此参数用户无需关注，设置为0即可。<br>注：预留的扩展参数，当前因后续架构升级，该参数已废弃，不对其进行业务处理。 \|	88	\| sid \| 此参数用户无需关注，设置为0即可。<br>注：预留的扩展参数，当前因后续架构升级，该参数已废弃，不对其进行业务处理。 \|
89	\| dstGap \| 相邻迭代间，目的操作数前一个分形结束地址与后一个分形起始地址的间隔，单位：512字节。取值范围：dstGap∈[0, 65535]。默认为0。<br>特性细节可参考：[非连续搬入](#非连续搬入)。<br>注：dstGap=0表示相邻repeat目的操作数起始地址间隔1个数据分形即连续存放。 \|	89	\| dstGap \| 相邻迭代间，目的操作数前一个分形结束地址与后一个分形起始地址的间隔，单位：512字节。取值范围：dstGap∈[0, 65535]。默认为0。<br>特性细节可参考：[非连续搬入](#非连续搬入)。<br>注：dstGap=0表示相邻repeat目的操作数起始地址间隔1个数据分形即连续存放。 \|
90	\| ifTranspose \| 是否启用转置功能，对每个分形矩阵进行转置，默认为false：<br>  • true：启用<br>  • false：不启用<br>特性细节可参考：[分形转置](#分形转置)。<br>注：只有L1 Buffer->L0A Buffer和L1 Buffer->L0B Buffer通路才能开启转置，开启转置功能时，源操作数、目的操作数仅支持b16数据类型。 \|	90	\| ifTranspose \| 是否启用转置功能，对每个分形矩阵进行转置，默认为false：<br>  • true：启用<br>  • false：不启用<br>特性细节可参考：[分形转置](#分形转置)。<br>注：只有L1 Buffer->L0A Buffer和L1 Buffer->L0B Buffer通路才能开启转置，开启转置功能时，源操作数、目的操作数仅支持b16数据类型。 \|
91	-\| addrMode \| 用于控制多次迭代场景下，源操作数中每一次迭代的分形矩阵索引ID是递增还是递减：<br>  • 0（默认）：递增，下一次repeat index = startIndex + srcStride×repeatTime；<br>  • 1：递减，下一次repeat index = startIndex - srcStride×repeatTime。<br>特性细节可参考：[控制地址更新方式](#控制地址更新方式)。<br>注：目前仅GM->L1 Buffer通路支持配置，保持默认值0即可，该参数不涉及性能。 \|	91	+\| addrMode \| 用于控制多次迭代场景下，源操作数中每一次迭代的分形矩阵索引ID是递增还是递减：<br>  • 0（默认）：递增，下一次repeat index = startIndex + srcStride \* repeatTimes；<br>  • 1：递减，下一次repeat index = startIndex - srcStride \* repeatTimes。<br>特性细节可参考：[控制地址更新方式](#控制地址更新方式)。<br>注：目前仅GM->L1 Buffer通路支持配置，保持默认值0即可，该参数不涉及性能。 \|
92		92
93	## 数据类型	93	## 数据类型
94		94

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬入/矩阵数据搬入至L0-Buffer/Load3D.md+1-1

	@@ -506,7 +506,7 @@ repeatMode =1的时候filter窗口中读取数据的位置保持不变，每个r
506		506
507	## 调用示例<a id="zh-cn_topic_0000002512171652_section088124295117"></a>	507	## 调用示例<a id="zh-cn_topic_0000002512171652_section088124295117"></a>
508		508
509	-如下示例中：A矩阵转置，shape为[k,m]，b32数据类型下，GM->L1 Buffer->L0A Buffer数据排布分别是ND、NZ和ZZ。但是L1 Buffer->L0A Buffer的时候，无法调用LoadDataWithTranspose指令对进行转置，因为在K轴方向两个连续的分形并不能合并为一个16\*16的方块，示意图如下：	509	+如下示例中：A矩阵转置，shape为[k,m]，b32数据类型下，GM->L1 Buffer->L0A Buffer数据排布分别是ND、NZ和ZZ。但是L1 Buffer->L0A Buffer的时候，无法调用LoadDataWithTranspose指令对其进行转置，因为在K轴方向两个连续的分形并不能合并为一个16\*16的方块，示意图如下：
510		510
511	![](../../../../../figures/load3d_l1tol0a_b32demo.png)	511	![](../../../../../figures/load3d_l1tol0a_b32demo.png)
512		512

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬入/矩阵数据搬入至L0-Buffer/LoadDataWithTranspose.md+12-12

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬入/矩阵数据搬入至L1-Buffer/GMToL1-Load2D指令搬运.md+1-1

	@@ -63,7 +63,7 @@ __aicore__ inline void LoadData(const LocalTensor<T>& dst, const GlobalTensor<T>
63	\| sid \| 此参数用户无需关注，设置为0即可。 \|	63	\| sid \| 此参数用户无需关注，设置为0即可。 \|
64	\| dstGap \| 相邻迭代间，目的操作数前一个分形结束地址与后一个分形起始地址的间隔，单位：512字节。取值范围：dstGap∈[0, 65535]。默认为0。<br>注：dstGap = 0表示相邻repeat目的操作数起始地址间隔1个数据分形即连续存放。 \|	64	\| dstGap \| 相邻迭代间，目的操作数前一个分形结束地址与后一个分形起始地址的间隔，单位：512字节。取值范围：dstGap∈[0, 65535]。默认为0。<br>注：dstGap = 0表示相邻repeat目的操作数起始地址间隔1个数据分形即连续存放。 \|
65	\| ifTranspose \| 是否启用转置功能，对每个分形矩阵进行转置，默认为false：<br>  • true：启用。<br>  • false：不启用。<br>注：本通路场景下不支持转置，参数无意义，保持默认即可。 \|	65	\| ifTranspose \| 是否启用转置功能，对每个分形矩阵进行转置，默认为false：<br>  • true：启用。<br>  • false：不启用。<br>注：本通路场景下不支持转置，参数无意义，保持默认即可。 \|
66	-\| addrMode \| 用于控制多次迭代场景下，源操作数中每一次迭代的分形矩阵索引ID是递增还是递减，0表示递增，下一次repeat index = startIndex + srcStride \* repeatTime。1表示递减，下一次repeat index = startIndex - srcStride \* repeatTime。 \|	66	+\| addrMode \| 用于控制多次迭代场景下，源操作数中每一次迭代的分形矩阵索引ID是递增还是递减，0表示递增，下一次repeat index = startIndex + srcStride \* repeatTimes。1表示递减，下一次repeat index = startIndex - srcStride \* repeatTimes。 \|
67		67
68	## 数据类型<a id="zh-cn_topic_0000002567745223_section4219135304818"></a>	68	## 数据类型<a id="zh-cn_topic_0000002567745223_section4219135304818"></a>
69		69

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/L0C到GM数据搬运（DataCopy）.md+2-2

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/L0C到GM数据搬运（Fixpipe）.md+3-3

	@@ -136,7 +136,7 @@
136	\| deqScalar \| 可选输入 \| scalar量化参数，表示单个scale值，quantPre量化模式为[随路量化](./关键特性说明/随路量化.md)时需要设置该参数。支持的数据类型为`uint64_t`。 \|	136	\| deqScalar \| 可选输入 \| scalar量化参数，表示单个scale值，quantPre量化模式为[随路量化](./关键特性说明/随路量化.md)时需要设置该参数。支持的数据类型为`uint64_t`。 \|
137	\| [reluEn](./关键特性说明/随路ReLU.md) \| 可选输入 \| 是否开启NormReLU的开关：<br>  • `false`：不开启NormReLU功能，默认为`false`；<br>  • `true`：开启NormReLU功能。 \|	137	\| [reluEn](./关键特性说明/随路ReLU.md) \| 可选输入 \| 是否开启NormReLU的开关：<br>  • `false`：不开启NormReLU功能，默认为`false`；<br>  • `true`：开启NormReLU功能。 \|
138	\| unitFlag \| 可选输入 \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag；<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位；<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>参数设置方案和特性细节可参考：Mmad计算中关键特性说明的[UnitFlag](../Mmad计算/关键特性说明/UnitFlag.md)章节。 \|	138	\| unitFlag \| 可选输入 \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag；<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位；<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>参数设置方案和特性细节可参考：Mmad计算中关键特性说明的[UnitFlag](../Mmad计算/关键特性说明/UnitFlag.md)章节。 \|
139	-\| params \| 可选输入 \| 用于选择和配置不同的随路格式转换（NZ2ZN/NZ2ND/NZ2DN），该参数为TransformParams类型的结构体。TransformParams结构体是一个基于模板参数的类型选择器，用于在编译时根据定义FixpipeParamsArch3510搬运参数时使用的模板参数，自动选择对应的参数类型。<br><pre>template \<CO2Layout format\><br>struct TransformParams {};<br>template \<\><br>struct TransformParams\<CO2Layout::NZ\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = uint8_t;<br>};<br>template \<\><br>struct TransformParams\<CO2Layout::ROW_MAJOR\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = Nz2NdParams;<br>};<br>template \<\><br>struct TransformParams\<CO2Layout::COLUMN_MAJOR\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = Nz2DnParams;<br>};</pre><br>CO2Layout布局类型为ROW_MAJOR时，该指令被定义为从L0C Buffer到目标位置的数据移动，并附带NZ2ND转换，结构体参数如下：<br><pre>struct Nz2NdParams {<br> uint16_t ndNum = 1;<br> uint16_t srcNdStride = 0;<br> uint32_t dstNdStride = 0; <br>};</pre>  • ndNum：源NZ矩阵的数目，也就是传输ND矩阵的数目，取值范围为ndNum∈[0, 65535]。注：ndNum=0表示不执行搬运，该接口将被视为NOP（空操作）。<br>  • srcNdStride：不同NZ矩阵起始地址之间的间隔，取值范围为srcNdStride∈[0, 65535]，单位为C0_SIZE。当ndNum配置为1时，srcNdStride配置为0即可，不生效。<br>  • dstNdStride：目的相邻ND矩阵起始地址之间的偏移，取值范围为dstNdStride∈[1, 2^32 -1]，单位为element。当ndNum配置为1时，dstNdStride配置为0即可，不生效。<br><br>CO2Layout布局类型为COLUMN_MAJOR时，该指令被定义为从L0C Buffer到目标位置的数据移动，并附带NZ2DN转换，结构体参数如下：<br><pre>struct Nz2DnParams {<br> uint16_t dnNum = 1;<br> uint16_t srcNzMatrixStride = 0;<br> uint32_t dstDnMatrixStride = 0;<br> uint16_t srcNzC0Stride = 0; <br>};</pre>  • dnNum：传输DN矩阵的数目，取值范围为dnNum∈[0, 65535]。注：dnNum=0表示不执行搬运，该接口将被视为NOP（空操作）。<br>  • srcNzMatrixStride：不同源NZ矩阵的偏移（头与头），取值范围srcNzMatrixStride∈[0, 65535]，单位C0_SIZE。当dnNum配置为1时，srcNzMatrixStride配置为0即可，不生效。<br>  • dstDnMatrixStride：目的相邻ND矩阵起始地址之间的偏移，取值范围dstDnMatrixStride∈[0, 2^32 -1]，单位element。当dnNum配置为1时，dstDnMatrixStride配置为0即可，不生效。<br>  • srcNzC0Stride：源矩阵NZ分形中相邻行的地址偏移（头与头），取值范围srcNzC0Stride∈[0, 65535]，单位C0_SIZE。当启用NZ2DN时，srcNzC0Stride不能为0。<br><br>CO2Layout布局类型为NZ时，为普通搬运DMA模式，表示从L0C Buffer到目标位置的正常数据移动。 \|	139	+\| params \| 可选输入 \| 用于选择和配置不同的随路格式转换（NZ2NZ/NZ2ND/NZ2DN），该参数为TransformParams类型的结构体。TransformParams结构体是一个基于模板参数的类型选择器，用于在编译时根据定义FixpipeParamsArch3510搬运参数时使用的模板参数，自动选择对应的参数类型。<br><pre>template \<CO2Layout format\><br>struct TransformParams {};<br>template \<\><br>struct TransformParams\<CO2Layout::NZ\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = uint8_t;<br>};<br>template \<\><br>struct TransformParams\<CO2Layout::ROW_MAJOR\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = Nz2NdParams;<br>};<br>template \<\><br>struct TransformParams\<CO2Layout::COLUMN_MAJOR\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = Nz2DnParams;<br>};</pre><br>CO2Layout布局类型为ROW_MAJOR时，该指令被定义为从L0C Buffer到目标位置的数据移动，并附带NZ2ND转换，结构体参数如下：<br><pre>struct Nz2NdParams {<br> uint16_t ndNum = 1;<br> uint16_t srcNdStride = 0;<br> uint32_t dstNdStride = 0; <br>};</pre>  • ndNum：源NZ矩阵的数目，也就是传输ND矩阵的数目，取值范围为ndNum∈[0, 65535]。注：ndNum=0表示不执行搬运，该接口将被视为NOP（空操作）。<br>  • srcNdStride：不同NZ矩阵起始地址之间的间隔，取值范围为srcNdStride∈[0, 65535]，单位为C0_SIZE。当ndNum配置为1时，srcNdStride配置为0即可，不生效。<br>  • dstNdStride：目的相邻ND矩阵起始地址之间的偏移，取值范围为dstNdStride∈[1, 2^32 -1]，单位为element。当ndNum配置为1时，dstNdStride配置为0即可，不生效。<br><br>CO2Layout布局类型为COLUMN_MAJOR时，该指令被定义为从L0C Buffer到目标位置的数据移动，并附带NZ2DN转换，结构体参数如下：<br><pre>struct Nz2DnParams {<br> uint16_t dnNum = 1;<br> uint16_t srcNzMatrixStride = 0;<br> uint32_t dstDnMatrixStride = 0;<br> uint16_t srcNzC0Stride = 0; <br>};</pre>  • dnNum：传输DN矩阵的数目，取值范围为dnNum∈[0, 65535]。注：dnNum=0表示不执行搬运，该接口将被视为NOP（空操作）。<br>  • srcNzMatrixStride：不同源NZ矩阵的偏移（头与头），取值范围srcNzMatrixStride∈[0, 65535]，单位C0_SIZE。当dnNum配置为1时，srcNzMatrixStride配置为0即可，不生效。<br>  • dstDnMatrixStride：目的相邻ND矩阵起始地址之间的偏移，取值范围dstDnMatrixStride∈[0, 2^32 -1]，单位element。当dnNum配置为1时，dstDnMatrixStride配置为0即可，不生效。<br>  • srcNzC0Stride：源矩阵NZ分形中相邻行的地址偏移（头与头），取值范围srcNzC0Stride∈[0, 65535]，单位C0_SIZE。当启用NZ2DN时，srcNzC0Stride不能为0。<br><br>CO2Layout布局类型为NZ时，为普通搬运DMA模式，表示从L0C Buffer到目标位置的正常数据移动。 \|
140	\| dualDstCtrl \| 可选输入 \| 此参数仅在L0C Buffer到Unified Buffer通路下有效。 \|	140	\| dualDstCtrl \| 可选输入 \| 此参数仅在L0C Buffer到Unified Buffer通路下有效。 \|
141	\| subBlockId \| 可选输入 \| 此参数仅在L0C Buffer到Unified Buffer通路下有效。 \|	141	\| subBlockId \| 可选输入 \| 此参数仅在L0C Buffer到Unified Buffer通路下有效。 \|
142	\| [isChannelSplit](./关键特性说明/F32-Channel-Split.md) \| 可选输入 \| 是否开启通道拆分的功能。默认为false，不开启该功能。仅在src和dst都为float时才能开启通道拆分，且不能同时开启ChannelSplit和NZ2ND/NZ2DN功能。 \|	142	\| [isChannelSplit](./关键特性说明/F32-Channel-Split.md) \| 可选输入 \| 是否开启通道拆分的功能。默认为false，不开启该功能。仅在src和dst都为float时才能开启通道拆分，且不能同时开启ChannelSplit和NZ2ND/NZ2DN功能。 \|
	@@ -169,7 +169,7 @@
169	\| quantPre \| 输入 \| QuantMode_t是一个枚举类型，用于控制量化模式，默认值为QuantMode_t::NoQuant，即不开启量化功能。QuantMode_t取值如下：<br>  • NoQuant，不开启量化功能；<br>  • F322F16，float cast成half，cast mode为CAST_RINT模式；<br>  • F322BF16，float cast成bfloat16_t，cast mode为CAST_RINT模式；<br>  • DEQF16，int32_t量化成half，scalar量化；<br>  • VDEQF16，int32_t量化成half，tensor量化；<br>  • QF322B8_PRE，float量化成uint8_t/int8_t，scalar量化；<br>  • VQF322B8_PRE，float量化成uint8_t/int8_t，tensor量化；<br>  • REQ8，int32_t量化成uint8_t/int8_t，scalar量化；<br>  • VREQ8，int32_t量化成uint8_t/int8_t，tensor量化。 \|	169	\| quantPre \| 输入 \| QuantMode_t是一个枚举类型，用于控制量化模式，默认值为QuantMode_t::NoQuant，即不开启量化功能。QuantMode_t取值如下：<br>  • NoQuant，不开启量化功能；<br>  • F322F16，float cast成half，cast mode为CAST_RINT模式；<br>  • F322BF16，float cast成bfloat16_t，cast mode为CAST_RINT模式；<br>  • DEQF16，int32_t量化成half，scalar量化；<br>  • VDEQF16，int32_t量化成half，tensor量化；<br>  • QF322B8_PRE，float量化成uint8_t/int8_t，scalar量化；<br>  • VQF322B8_PRE，float量化成uint8_t/int8_t，tensor量化；<br>  • REQ8，int32_t量化成uint8_t/int8_t，scalar量化；<br>  • VREQ8，int32_t量化成uint8_t/int8_t，tensor量化。 \|
170	\| deqScalar \| 输入 \| scalar量化参数，表示单个scale值，quantPre量化模式为scalar量化时需要设置该参数。支持的数据类型为uint64_t。 \|	170	\| deqScalar \| 输入 \| scalar量化参数，表示单个scale值，quantPre量化模式为scalar量化时需要设置该参数。支持的数据类型为uint64_t。 \|
171	\| ndNum \| 输入 \| 源NZ矩阵的数目，也就是传输ND矩阵的数目，取值范围为ndNum∈[1, 65535]。 \|	171	\| ndNum \| 输入 \| 源NZ矩阵的数目，也就是传输ND矩阵的数目，取值范围为ndNum∈[1, 65535]。 \|
172	-\| srcNdStride \| 输入 \| 不同NZ矩阵起始地址之间的间隔，取值范围为srcNdStride∈[1, 512]，单位为数据块（16 \* C0_Size）。当ndNum配置为1时，srcNdStride配置为0即可，不生效。 \|	172	+\| srcNdStride \| 输入 \| 不同NZ矩阵起始地址之间的间隔。当ndNum大于1时，取值范围为srcNdStride∈[1, 512]；当ndNum配置为1时，srcNdStride可配置为0，此时不生效。单位为数据块（16 \* C0_Size）。 \|
173	\| dstNdStride \| 输入 \| 目的相邻ND矩阵起始地址之间的偏移，取值范围为dstNdStride∈[1, 65535]，单位为element。当ndNum配置为1时，dstNdStride配置为0即可，不生效。 \|	173	\| dstNdStride \| 输入 \| 目的相邻ND矩阵起始地址之间的偏移，取值范围为dstNdStride∈[1, 65535]，单位为element。当ndNum配置为1时，dstNdStride配置为0即可，不生效。 \|
174	\| reluEn \| 输入 \| 是否开启ReLU的开关：<br>  • false：不开启ReLU功能；<br>  • true：开启ReLU功能。 \|	174	\| reluEn \| 输入 \| 是否开启ReLU的开关：<br>  • false：不开启ReLU功能；<br>  • true：开启ReLU功能。 \|
175	\| unitFlag \| 输入 \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出，该功能不适用于在L0C Buffer累加的场景。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag；<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不会设置寄存器；<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，会将unitFlag关闭。 \|	175	\| unitFlag \| 输入 \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出，该功能不适用于在L0C Buffer累加的场景。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag；<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不会设置寄存器；<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，会将unitFlag关闭。 \|
	@@ -249,7 +249,7 @@
249	- 目标数据不能有重叠。如果对目的地址有重叠写入，硬件不会报告任何警告和错误，也不保证重叠数据的写入顺序。	249	- 目标数据不能有重叠。如果对目的地址有重叠写入，硬件不会报告任何警告和错误，也不保证重叠数据的写入顺序。
250	- unitFlag特性开启需要配合Mmad同时开启。	250	- unitFlag特性开启需要配合Mmad同时开启。
251	<!-- npu="950" id28 -->	251	<!-- npu="950" id28 -->
252	-- 当启用NZ2DN且srcNzC0Stride不等于1时，不能同时开启unitflag。	252	+- 当启用NZ2DN且srcNzC0Stride不等于1时，不能同时开启unitFlag。
253	<!-- end id28 -->	253	<!-- end id28 -->
254	<!-- npu="A3,910b" id27 -->	254	<!-- npu="A3,910b" id27 -->
255	- 针对如下产品型号，特殊值/边界值约束说明如下：	255	- 针对如下产品型号，特殊值/边界值约束说明如下：

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/L0C到L1数据搬运（DataCopy）.md+1-1

	@@ -98,7 +98,7 @@ DataCopy矩阵搬出接口支持多种随路能力的组合，需要设置不同
98	\| dstStride \| 目的NZ矩阵中相邻Z排布的起始地址偏移，取值不为0，单位为datablock（32字节）。 \|	98	\| dstStride \| 目的NZ矩阵中相邻Z排布的起始地址偏移，取值不为0，单位为datablock（32字节）。 \|
99	\| srcStride \| 源NZ矩阵中相邻Z排布的起始地址偏移，取值范围为srcStride∈[0, 65535]，单位为C0_Size（16*sizeof(T)），T为src的数据类型，其值应填成mSize对16向上取整。 \|	99	\| srcStride \| 源NZ矩阵中相邻Z排布的起始地址偏移，取值范围为srcStride∈[0, 65535]，单位为C0_Size（16*sizeof(T)），T为src的数据类型，其值应填成mSize对16向上取整。 \|
100	\| unitFlag \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag。<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位。<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>参数设置方案和特性细节可参考：Mmad计算中关键特性说明的[UnitFlag](../Mmad计算/关键特性说明/UnitFlag.md)章节。 \|	100	\| unitFlag \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag。<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位。<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>参数设置方案和特性细节可参考：Mmad计算中关键特性说明的[UnitFlag](../Mmad计算/关键特性说明/UnitFlag.md)章节。 \|
101	-\| clipReluPre \| <!-- npu="950,A3,910b" id13 -->该参数仅在Atlas 200I/500 A2 推理产品支持。<br><!-- end id13 --><!-- npu="310b" id14 -->用于配置是否开启ClipReLU操作，参数类型为uint8_t，取值如下：0，不开启ClipReLU；1，开启ClipReLU，此时需要调用[SetFixPipeClipRelu](./寄存器配置说明/SetFixPipeClipRelu.md)来设置clipReLU的最大值。<br>  •该操作在随路量化后进行，quantPre配置后才能使用，当前支持的量化模式有F322F16/DEQF16/VDEQF16/QF322B8_PRE/VQF322B8_PRE/REQ8/VREQ8。<br><!-- end id14 --> \|	101	+\| clipReluPre \| <!-- npu="950,A3,910b" id13 -->该参数仅在Atlas 200I/500 A2 推理产品支持。<br><!-- end id13 --><!-- npu="310b" id14 -->用于配置是否开启ClipReLU操作，参数类型为uint8_t，取值如下：0，不开启ClipReLU；1，开启ClipReLU，此时需要调用[SetFixPipeClipRelu](./寄存器配置说明/SetFixPipeClipRelu.md)来设置ClipReLU的最大值。<br>  •该操作在随路量化后进行，quantPre配置后才能使用，当前支持的量化模式有F322F16/DEQF16/VDEQF16/QF322B8_PRE/VQF322B8_PRE/REQ8/VREQ8。<br><!-- end id14 --> \|
102	\| eltWiseOp \| <!-- npu="950,A3,910b" id15 -->该参数仅在Atlas 200I/500 A2 推理产品支持。<br><!-- end id15 --><!-- npu="310b" id16 -->用于配置是否开启Elementwise操作及操作模式。Elementwise操作是指进行随路量化后，可以逐个元素加/减一个LocalTensor，大小为mSize * nSize，具体LocalTensor地址相关参数需要调用[SetFixPipeAddr](./寄存器配置说明/SetFixPipeAddr.md)来设置。<br>eltWiseOp参数类型为uint8_t，取值如下：<br>  • 0：不开启Elementwise；<br>  • 1：Elementwise Addition；<br>  • 2：Elementwise Subtraction。<br><!-- end id16 --> \|	102	\| eltWiseOp \| <!-- npu="950,A3,910b" id15 -->该参数仅在Atlas 200I/500 A2 推理产品支持。<br><!-- end id15 --><!-- npu="310b" id16 -->用于配置是否开启Elementwise操作及操作模式。Elementwise操作是指进行随路量化后，可以逐个元素加/减一个LocalTensor，大小为mSize * nSize，具体LocalTensor地址相关参数需要调用[SetFixPipeAddr](./寄存器配置说明/SetFixPipeAddr.md)来设置。<br>eltWiseOp参数类型为uint8_t，取值如下：<br>  • 0：不开启Elementwise；<br>  • 1：Elementwise Addition；<br>  • 2：Elementwise Subtraction。<br><!-- end id16 --> \|
103	\| quantPre \| 用于控制量化模式，QuantMode_t类型，具体定义如下：<br>  • half/bfloat16_t输出，此参数需配置为QuantMode_t::F322F16/QuantMode_t::F322BF16。<br>  •配置为scalar量化时，需要调用[SetFixpipePreQuantFlag](./寄存器配置说明/SetFixpipePreQuantFlag.md)接口来设置scalar量化参数。<br>  •配置为tensor量化时，需要调用[SetFixPipeConfig](./寄存器配置说明/SetFixPipeConfig.md)来设置tensor量化参数，其中tensor量化参数需要通过DataCopy从L1 Buffer搬运至Fixpipe Buffer。<br>注：此参数需要用户手动配置，不会自动推导配置对应量化模式。此通路不支持NoQuant模式。<br><pre>enum QuantMode_t<br>{<br> NoQuant, // 不开启量化功能，此通路不支持<br> F322F16, // Float32_2_Float16：float cast成half，cast mode为CAST_RINT模式<br> F322BF16, // Float32_2_BFloat16：float cast成bfloat16_t，cast mode为CAST_RINT模式<br> DEQF16, // DeQuant_Float16：int32_t量化成half，scalar量化<br> VDEQF16, // Vector_DeQuant_Float16：int32_t量化成half，tensor量化<br> QF322B8_PRE, // Quant_Float32_2_B8：float量化成int8_t/uint8_t，scalar量化<br> VQF322B8_PRE, // Vector_Quant_Float32_2_B8：float量化成int8_t/uint8_t，tensor量化<br> REQ8, // ReQuant_int8：int32_t量化成int8_t/uint8_t，scalar量化<br> VREQ8, // Vector_ReQuant_int8：int32_t量化成int8_t/uint8_t，tensor量化<br>};</pre> \|	103	\| quantPre \| 用于控制量化模式，QuantMode_t类型，具体定义如下：<br>  • half/bfloat16_t输出，此参数需配置为QuantMode_t::F322F16/QuantMode_t::F322BF16。<br>  •配置为scalar量化时，需要调用[SetFixpipePreQuantFlag](./寄存器配置说明/SetFixpipePreQuantFlag.md)接口来设置scalar量化参数。<br>  •配置为tensor量化时，需要调用[SetFixPipeConfig](./寄存器配置说明/SetFixPipeConfig.md)来设置tensor量化参数，其中tensor量化参数需要通过DataCopy从L1 Buffer搬运至Fixpipe Buffer。<br>注：此参数需要用户手动配置，不会自动推导配置对应量化模式。此通路不支持NoQuant模式。<br><pre>enum QuantMode_t<br>{<br> NoQuant, // 不开启量化功能，此通路不支持<br> F322F16, // Float32_2_Float16：float cast成half，cast mode为CAST_RINT模式<br> F322BF16, // Float32_2_BFloat16：float cast成bfloat16_t，cast mode为CAST_RINT模式<br> DEQF16, // DeQuant_Float16：int32_t量化成half，scalar量化<br> VDEQF16, // Vector_DeQuant_Float16：int32_t量化成half，tensor量化<br> QF322B8_PRE, // Quant_Float32_2_B8：float量化成int8_t/uint8_t，scalar量化<br> VQF322B8_PRE, // Vector_Quant_Float32_2_B8：float量化成int8_t/uint8_t，tensor量化<br> REQ8, // ReQuant_int8：int32_t量化成int8_t/uint8_t，scalar量化<br> VREQ8, // Vector_ReQuant_int8：int32_t量化成int8_t/uint8_t，tensor量化<br>};</pre> \|
104	\| reluPre \| 用于配置ReLU操作的模式，类型为uint8_t，取值如下：<br>  • 0：不开启ReLU<br>  • 1：Normal ReLU \|	104	\| reluPre \| 用于配置ReLU操作的模式，类型为uint8_t，取值如下：<br>  • 0：不开启ReLU<br>  • 1：Normal ReLU \|

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/L0C到L1数据搬运（Fixpipe）.md+2-2

	@@ -120,7 +120,7 @@
120	\| deqScalar \| 可选输入 \| scalar量化参数，表示单个scale值，quantPre量化模式为[随路量化](./关键特性说明/随路量化.md)时需要设置该参数。支持的数据类型为`uint64_t`。 \|	120	\| deqScalar \| 可选输入 \| scalar量化参数，表示单个scale值，quantPre量化模式为[随路量化](./关键特性说明/随路量化.md)时需要设置该参数。支持的数据类型为`uint64_t`。 \|
121	\| [reluEn](./关键特性说明/随路ReLU.md) \| 可选输入 \| 是否开启NormReLU的开关：<br>  • `false`：不开启NormReLU功能，默认为`false`；<br>  • `true`：开启NormReLU功能。 \|	121	\| [reluEn](./关键特性说明/随路ReLU.md) \| 可选输入 \| 是否开启NormReLU的开关：<br>  • `false`：不开启NormReLU功能，默认为`false`；<br>  • `true`：开启NormReLU功能。 \|
122	\| unitFlag \| 可选输入 \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag；<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位；<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>参数设置方案和特性细节可参考：Mmad计算中关键特性说明的[UnitFlag](../Mmad计算/关键特性说明/UnitFlag.md)章节。 \|	122	\| unitFlag \| 可选输入 \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag；<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位；<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>参数设置方案和特性细节可参考：Mmad计算中关键特性说明的[UnitFlag](../Mmad计算/关键特性说明/UnitFlag.md)章节。 \|
123	-\| params \| 可选输入 \| 用于选择和配置不同的随路格式转换（NZ2ZN/NZ2ND/NZ2DN），该参数为TransformParams类型的结构体。TransformParams结构体是一个基于模板参数的类型选择器，用于在编译时根据定义FixpipeParamsArch3510搬运参数时使用的模板参数，自动选择对应的参数类型。<br><pre>template \<CO2Layout format\><br>struct TransformParams {};<br>template \<\><br>struct TransformParams\<CO2Layout::NZ\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = uint8_t;<br>};<br>template \<\><br>struct TransformParams\<CO2Layout::ROW_MAJOR\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = Nz2NdParams;<br>};<br>template \<\><br>struct TransformParams\<CO2Layout::COLUMN_MAJOR\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = Nz2DnParams;<br>};</pre><br>CO2Layout布局类型为ROW_MAJOR时，该指令被定义为从L0C Buffer到目标位置的数据移动，并附带NZ2ND转换，结构体参数如下：<br><pre>struct Nz2NdParams {<br> uint16_t ndNum = 1;<br> uint16_t srcNdStride = 0;<br> uint32_t dstNdStride = 0; <br>};</pre>  • ndNum：源NZ矩阵的数目，也就是传输ND矩阵的数目，取值范围为ndNum∈[0, 65535]。注：ndNum=0表示不执行搬运，该接口将被视为NOP（空操作）。<br>  • srcNdStride：不同NZ矩阵起始地址之间的间隔，取值范围为srcNdStride∈[0, 65535]，单位为C0_SIZE。当ndNum配置为1时，srcNdStride配置为0即可，不生效。<br>  • dstNdStride：目的相邻ND矩阵起始地址之间的偏移，取值范围为dstNdStride∈[1, 2^32 -1]，单位为element。当ndNum配置为1时，dstNdStride配置为0即可，不生效。<br><br>CO2Layout布局类型为COLUMN_MAJOR时，该指令被定义为从L0C Buffer到目标位置的数据移动，并附带NZ2DN转换，结构体参数如下：<br><pre>struct Nz2DnParams {<br> uint16_t dnNum = 1;<br> uint16_t srcNzMatrixStride = 0;<br> uint32_t dstDnMatrixStride = 0;<br> uint16_t srcNzC0Stride = 0; <br>};</pre>  • dnNum：传输DN矩阵的数目，取值范围为dnNum∈[0, 65535]。注：dnNum=0表示不执行搬运，该接口将被视为NOP（空操作）。<br>  • srcNzMatrixStride：不同源NZ矩阵的偏移（头与头），取值范围srcNzMatrixStride∈[0, 65535]，单位C0_SIZE。当dnNum配置为1时，srcNzMatrixStride配置为0即可，不生效。<br>  • dstDnMatrixStride：目的相邻ND矩阵起始地址之间的偏移，取值范围dstDnMatrixStride∈[0, 2^32 -1]，单位element。当dnNum配置为1时，dstDnMatrixStride配置为0即可，不生效。<br>  • srcNzC0Stride：源矩阵NZ分形中相邻行的地址偏移（头与头），取值范围srcNzC0Stride∈[0, 65535]，单位C0_SIZE。当启用NZ2DN时，srcNzC0Stride不能为0。<br><br>CO2Layout布局类型为NZ时，为普通搬运DMA模式，表示从L0C Buffer到目标位置的正常数据移动。 \|	123	+\| params \| 可选输入 \| 用于选择和配置不同的随路格式转换（NZ2NZ/NZ2ND/NZ2DN），该参数为TransformParams类型的结构体。TransformParams结构体是一个基于模板参数的类型选择器，用于在编译时根据定义FixpipeParamsArch3510搬运参数时使用的模板参数，自动选择对应的参数类型。<br><pre>template \<CO2Layout format\><br>struct TransformParams {};<br>template \<\><br>struct TransformParams\<CO2Layout::NZ\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = uint8_t;<br>};<br>template \<\><br>struct TransformParams\<CO2Layout::ROW_MAJOR\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = Nz2NdParams;<br>};<br>template \<\><br>struct TransformParams\<CO2Layout::COLUMN_MAJOR\> {<br> \_\_aicore\_\_ inline TransformParams(){};<br> using PARAMS = Nz2DnParams;<br>};</pre><br>CO2Layout布局类型为ROW_MAJOR时，该指令被定义为从L0C Buffer到目标位置的数据移动，并附带NZ2ND转换，结构体参数如下：<br><pre>struct Nz2NdParams {<br> uint16_t ndNum = 1;<br> uint16_t srcNdStride = 0;<br> uint32_t dstNdStride = 0; <br>};</pre>  • ndNum：源NZ矩阵的数目，也就是传输ND矩阵的数目，取值范围为ndNum∈[0, 65535]。注：ndNum=0表示不执行搬运，该接口将被视为NOP（空操作）。<br>  • srcNdStride：不同NZ矩阵起始地址之间的间隔，取值范围为srcNdStride∈[0, 65535]，单位为C0_SIZE。当ndNum配置为1时，srcNdStride配置为0即可，不生效。<br>  • dstNdStride：目的相邻ND矩阵起始地址之间的偏移，取值范围为dstNdStride∈[1, 2^32 -1]，单位为element。当ndNum配置为1时，dstNdStride配置为0即可，不生效。<br><br>CO2Layout布局类型为COLUMN_MAJOR时，该指令被定义为从L0C Buffer到目标位置的数据移动，并附带NZ2DN转换，结构体参数如下：<br><pre>struct Nz2DnParams {<br> uint16_t dnNum = 1;<br> uint16_t srcNzMatrixStride = 0;<br> uint32_t dstDnMatrixStride = 0;<br> uint16_t srcNzC0Stride = 0; <br>};</pre>  • dnNum：传输DN矩阵的数目，取值范围为dnNum∈[0, 65535]。注：dnNum=0表示不执行搬运，该接口将被视为NOP（空操作）。<br>  • srcNzMatrixStride：不同源NZ矩阵的偏移（头与头），取值范围srcNzMatrixStride∈[0, 65535]，单位C0_SIZE。当dnNum配置为1时，srcNzMatrixStride配置为0即可，不生效。<br>  • dstDnMatrixStride：目的相邻ND矩阵起始地址之间的偏移，取值范围dstDnMatrixStride∈[0, 2^32 -1]，单位element。当dnNum配置为1时，dstDnMatrixStride配置为0即可，不生效。<br>  • srcNzC0Stride：源矩阵NZ分形中相邻行的地址偏移（头与头），取值范围srcNzC0Stride∈[0, 65535]，单位C0_SIZE。当启用NZ2DN时，srcNzC0Stride不能为0。<br><br>CO2Layout布局类型为NZ时，为普通搬运DMA模式，表示从L0C Buffer到目标位置的正常数据移动。 \|
124	\| dualDstCtrl \| 可选输入 \| 此参数仅在L0C Buffer到Unified Buffer通路下有效。 \|	124	\| dualDstCtrl \| 可选输入 \| 此参数仅在L0C Buffer到Unified Buffer通路下有效。 \|
125	\| subBlockId \| 可选输入 \| 此参数仅在L0C Buffer到Unified Buffer通路下有效。 \|	125	\| subBlockId \| 可选输入 \| 此参数仅在L0C Buffer到Unified Buffer通路下有效。 \|
126	\| [isChannelSplit](./关键特性说明/F32-Channel-Split.md) \| 可选输入 \| 此参数仅在L0C Buffer到Global Memory通路下有效。 \|	126	\| [isChannelSplit](./关键特性说明/F32-Channel-Split.md) \| 可选输入 \| 此参数仅在L0C Buffer到Global Memory通路下有效。 \|
	@@ -206,7 +206,7 @@
206	- 目标数据不能有重叠。如果对目的地址有重叠写入，硬件不会报告任何警告和错误，也不保证重叠数据的写入顺序。	206	- 目标数据不能有重叠。如果对目的地址有重叠写入，硬件不会报告任何警告和错误，也不保证重叠数据的写入顺序。
207	- unitFlag特性开启需要配合Mmad同时开启。	207	- unitFlag特性开启需要配合Mmad同时开启。
208	<!-- npu="950" id27 -->	208	<!-- npu="950" id27 -->
209	-- 当启用NZ2DN且srcNzC0Stride不等于1时，不能同时开启unitflag。	209	+- 当启用NZ2DN且srcNzC0Stride不等于1时，不能同时开启unitFlag。
210	<!-- end id27 -->	210	<!-- end id27 -->
211	<!-- npu="A3,910b" id28 -->	211	<!-- npu="A3,910b" id28 -->
212	- 针对如下产品型号，特殊值/边界值约束说明如下：	212	- 针对如下产品型号，特殊值/边界值约束说明如下：

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/L0C到UB数据搬运（Fixpipe）.md+4-4

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/关键特性说明/F32-Channel-Split.md+1-1

	@@ -25,7 +25,7 @@
25	<!-- end id1 -->	25	<!-- end id1 -->
26		26
27	<!-- npu="950" id4 -->	27	<!-- npu="950" id4 -->
28	-- 针对Ascend 950PR/Ascend 950DT，当Fixpipe指令的输入和输出数据类型都为float,NZ输出到GM，并且不开启NZ2ND、NZ2DN与unitflag时，才能够开启Channel Split功能。当前搬出到L1 Buffer和UB不支持该功能。	28	+- 针对Ascend 950PR/Ascend 950DT，当Fixpipe指令的输入和输出数据类型都为float,NZ输出到GM，并且不开启NZ2ND、NZ2DN与unitFlag时，才能够开启Channel Split功能。当前搬出到L1 Buffer和UB不支持该功能。
29	<!-- end id4 -->	29	<!-- end id4 -->
30		30
31	## 使用示例	31	## 使用示例

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/寄存器配置说明/SetFixPipeAddr.md+2-2

	@@ -61,7 +61,7 @@ LocalTensor支持的数据类型为：half。
61	DataCopy随路量化搬运后，可以逐个元素加/减一个大小为mSize \* nSize的LocalTensor，具体LocalTensor地址相关参数需要调用`SetFixPipeAddr`来设置。	61	DataCopy随路量化搬运后，可以逐个元素加/减一个大小为mSize \* nSize的LocalTensor，具体LocalTensor地址相关参数需要调用`SetFixPipeAddr`来设置。
62		62
63	```cpp	63	```cpp
64	-__aicore__ inline void SetEleSrcPara(const LocalTensor <half>& eleWiseData, uint16_t c0ChStride)	64	+__aicore__ inline void SetEleSrcPara(const LocalTensor<half>& eleWiseData, uint16_t c0ChStride)
65	{	65	{
66	AscendC::SetFixPipeAddr(eleWiseData, c0ChStride);	66	AscendC::SetFixPipeAddr(eleWiseData, c0ChStride);
67	}	67	}
	@@ -288,4 +288,4 @@ __aicore__ inline void SetEleSrcPara(const LocalTensor <half>& eleWiseData, uint
288	}	288	}
289	KERNEL_CUBE_DATACOPY(half, int8_t, int8_t, int32_t, 128, 1, 1, QuantMode_t::DEQF16);	289	KERNEL_CUBE_DATACOPY(half, int8_t, int8_t, int32_t, 128, 1, 1, QuantMode_t::DEQF16);
290	```	290	```
291	-	291	+

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/寄存器配置说明/SetFixPipeClipRelu.md+3-3

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/寄存器配置说明/SetFixPipeConfig.md+2-2

	@@ -60,7 +60,7 @@ __aicore__ inline void SetFixPipeConfig(const LocalTensor<T>& preData, bool isUn
60	\| reluPre \| 输入 \| 源操作数，ReLU操作时参与计算的tensor，类型为LocalTensor，支持的物理内存为Fixpipe Buffer（TPosition为C2PIPE2GM），为后续的功能扩展做保留，传入一个空LocalTensor即可。 \|	60	\| reluPre \| 输入 \| 源操作数，ReLU操作时参与计算的tensor，类型为LocalTensor，支持的物理内存为Fixpipe Buffer（TPosition为C2PIPE2GM），为后续的功能扩展做保留，传入一个空LocalTensor即可。 \|
61	\| quantPre \| 输入 \| 源操作数，quant tensor，量化操作时参与计算的tensor，类型为LocalTensor，支持的物理内存为Fixpipe Buffer（TPosition为C2PIPE2GM）。 \|	61	\| quantPre \| 输入 \| 源操作数，quant tensor，量化操作时参与计算的tensor，类型为LocalTensor，支持的物理内存为Fixpipe Buffer（TPosition为C2PIPE2GM）。 \|
62	\| isUnitFlag \| 输入 \| unitFlag清空标志位，类型为bool。如果设为true，L0C Buffer中所有内存块的单元标志位都会设为0。旨在消除之前操作中unitFlag产生的多余影响。 \|	62	\| isUnitFlag \| 输入 \| unitFlag清空标志位，类型为bool。如果设为true，L0C Buffer中所有内存块的单元标志位都会设为0。旨在消除之前操作中unitFlag产生的多余影响。 \|
63	-\| preData \| 输入 \| 支持设置一个Tensor，通过开关控制是ReLU Tensor还是Quant Tensor，支持的物理内存为Fixpipe Buffer（TPosition为C2PIPE2GM）。当前仅支持传入quant Tensor。对于单独设置unitFlag清空标志位isUnitFlag的场景，支持的物理地址为L0C Buffer（TPosition设置为CO1）。 \|	63	+\| preData \| 输入 \| 支持设置一个tensor，通过开关控制是ReLU tensor还是Quant tensor，支持的物理内存为Fixpipe Buffer（TPosition为C2PIPE2GM）。当前仅支持传入quant tensor。对于单独设置unitFlag清空标志位isUnitFlag的场景，支持的物理地址为L0C Buffer（TPosition设置为CO1）。 \|
64		64
65	## 数据类型	65	## 数据类型
66		66
	@@ -72,7 +72,7 @@ __aicore__ inline void SetFixPipeConfig(const LocalTensor<T>& preData, bool isUn
72		72
73	## 约束说明<a name="zh-cn_topic_0000002547729135_section633mcpsimp"></a>	73	## 约束说明<a name="zh-cn_topic_0000002547729135_section633mcpsimp"></a>
74		74
75	-quantPre和reluPre必须是Fixpipe Buffer上的Tensor。	75	+quantPre和reluPre必须是Fixpipe Buffer上的tensor。
76		76
77	## 调用示例<a name="section642mcpsimp"></a>	77	## 调用示例<a name="section642mcpsimp"></a>
78		78

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/寄存器配置说明/SetFixpipeNz2ndFlag.md+1-1

	@@ -65,7 +65,7 @@ Atlas 200I/500 A2 推理产品
65	\| 参数名称 \| 输入/输出 \| 含义 \|	65	\| 参数名称 \| 输入/输出 \| 含义 \|
66	\| ---------- \| ---------- \| ---------- \|	66	\| ---------- \| ---------- \| ---------- \|
67	\| ndNum \| 输入 \| nd矩阵的数量，取值范围∈[1, 65535]。（[具体含义图示参见NZ2ND介绍特性章节](../关键特性说明/NZ2ND.md)） \|	67	\| ndNum \| 输入 \| nd矩阵的数量，取值范围∈[1, 65535]。（[具体含义图示参见NZ2ND介绍特性章节](../关键特性说明/NZ2ND.md)） \|
68	-\| srcNdStride \| 输入 \| 以分形大小为单位的源步长，源相邻nz矩阵的偏移（头与头）。当ndNum配置为1时，srcNdStride配置为0即可，不生效。取值范围∈[1, 512]，单位为1024B。 \|	68	+\| srcNdStride \| 输入 \| 以分形大小为单位的源步长，源相邻nz矩阵的偏移（头与头）。当ndNum大于1时，取值范围∈[1, 512]；当ndNum配置为1时，srcNdStride可配置为0，此时不生效。单位为1024B。 \|
69	\| dstNdStride \| 输入 \| 目的相邻nd矩阵的偏移（头与头）。取值范围dstNdStride∈[1, 65535]，单位为element元素。 \|	69	\| dstNdStride \| 输入 \| 目的相邻nd矩阵的偏移（头与头）。取值范围dstNdStride∈[1, 65535]，单位为element元素。 \|
70	<!-- end id11 -->	70	<!-- end id11 -->
71		71

Mdocs/api/SIMD-API/基础API/矩阵计算（ISASI）/矩阵计算的搬出/寄存器配置说明/SetFixpipePreQuantFlag.md+0-1

	@@ -37,7 +37,6 @@ DataCopy数据搬运（[L0C Buffer-\>GM](../L0C到GM数据搬运（DataCopy）.m
37	## 函数原型<a name="zh-cn_topic_0000002516049318_section620mcpsimp"></a>	37	## 函数原型<a name="zh-cn_topic_0000002516049318_section620mcpsimp"></a>
38		38
39	```cpp	39	```cpp
40	-template<typename T>
41	__aicore__ inline void SetFixpipePreQuantFlag(uint64_t config)	40	__aicore__ inline void SetFixpipePreQuantFlag(uint64_t config)
42	```	41	```
43		42

Mdocs/api/SIMD-API/高阶API/张量变换/GetBroadCastMaxMinTmpSize.md+1-1

	@@ -2,7 +2,7 @@
2		2
3	## 功能说明	3	## 功能说明
4		4
5	-kernel侧Broadcast接口的计算需要开发者预留/申请临时空间，本接口用于在host侧获取预留/申请的最大最小临时空间大小，开发者基于此范围选择合适的空间大小作为Tiling参数传递到kernel侧使用。	5	+kernel侧Broadcast接口的计算需要开发者预留/申请临时空间，本接口用于在host侧获取预留/申请的最大和最小临时空间大小，开发者基于此范围选择合适的空间大小作为Tiling参数传递到kernel侧使用。
6		6
7	- 为保证功能正确，预留/申请的临时空间大小不能小于最小临时空间大小；	7	- 为保证功能正确，预留/申请的临时空间大小不能小于最小临时空间大小；
8	- 在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。	8	- 在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。

Mdocs/api/SIMD-API/高阶API/张量变换/GetTransDataMaxMinTmpSize.md+1-1

	@@ -2,7 +2,7 @@
2		2
3	## 功能说明	3	## 功能说明
4		4
5	-kernel侧TransData接口的计算需要开发者预留/申请临时空间，本接口用于在host侧获取预留/申请的最大最小临时空间大小，开发者基于此范围选择合适的空间大小作为Tiling参数传递到kernel侧使用。	5	+kernel侧TransData接口的计算需要开发者预留/申请临时空间，本接口用于在host侧获取预留/申请的最大和最小临时空间大小，开发者基于此范围选择合适的空间大小作为Tiling参数传递到kernel侧使用。
6		6
7	- 为保证功能正确，预留/申请的临时空间大小不能小于最小临时空间大小。	7	- 为保证功能正确，预留/申请的临时空间大小不能小于最小临时空间大小。
8	- 在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。该接口最大临时空间当前等于最小临时空间。	8	- 在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。该接口最大临时空间当前等于最小临时空间。

Mdocs/api/SIMD-API/高阶API/张量变换/Pad-Tiling.md+5-5

	@@ -26,8 +26,8 @@ void PadTilingFunc(const ge::Shape srcShape, const ge::Shape oriSrcShape, const
26	\| --- \| --- \| --- \|	26	\| --- \| --- \| --- \|
27	\| srcShape \| 输入 \| 输入Tensor的shape信息，shape为二维。 \|	27	\| srcShape \| 输入 \| 输入Tensor的shape信息，shape为二维。 \|
28	\| typeSize \| 输入 \| 输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。 \|	28	\| typeSize \| 输入 \| 输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。 \|
29	-\| maxValue \| 输出 \| Pad接口能完成计算所需最大临时空间大小。<br><br>Pad接口能完成计算所需的最大临时空间大小，超出该值的空间不会被该接口使用。在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。 <br>maxValue仅作为参考值，有可能大于Unified Buffer剩余空间的大小，该场景下，开发者需要根据Unified Buffer剩余空间的大小来选取合适的临时空间大小。 \|	29	+\| maxValue \| 输出 \| Pad接口能完成计算所需的最大临时空间大小，超出该值的空间不会被该接口使用。在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。<br>maxValue仅作为参考值，有可能大于Unified Buffer剩余空间的大小，该场景下，开发者需要根据Unified Buffer剩余空间的大小来选取合适的临时空间大小。 \|
30	-\| minValue \| 输出 \| Pad接口能完成计算所需最小临时空间大小。<br><br>Pad接口能完成计算所需最小临时空间大小。为保证功能正确，接口计算时预留/申请的临时空间不能小于该数值。 \|	30	+\| minValue \| 输出 \| Pad接口能完成计算所需最小临时空间大小。为保证功能正确，接口计算时预留/申请的临时空间不能小于该数值。 \|
31		31
32	表2 PadTilingFunc接口参数说明	32	表2 PadTilingFunc接口参数说明
33		33
	@@ -35,7 +35,7 @@ void PadTilingFunc(const ge::Shape srcShape, const ge::Shape oriSrcShape, const
35	\| --- \| --- \| --- \|	35	\| --- \| --- \| --- \|
36	\| srcShape \| 输入 \| 输入Tensor的shape信息，shape为二维。（有效数据+冗余数据） \|	36	\| srcShape \| 输入 \| 输入Tensor的shape信息，shape为二维。（有效数据+冗余数据） \|
37	\| oriSrcShape \| 输入 \| 输入Tensor的原始shape信息，shape为二维。（有效数据） \|	37	\| oriSrcShape \| 输入 \| 输入Tensor的原始shape信息，shape为二维。（有效数据） \|
38	-\| stackBufferSize \| 输入 \| 可供Pad接口计算的空间大小。 \|	38	+\| stackBufferSize \| 输入 \| 可供Pad接口计算的临时空间大小，单位为Byte。 \|
39	\| typeSize \| 输入 \| 输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。 \|	39	\| typeSize \| 输入 \| 输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。 \|
40	\| tiling \| 输出 \| 输出Pad接口所需的tiling信息。 \|	40	\| tiling \| 输出 \| 输出Pad接口所需的tiling信息。 \|
41		41
	@@ -62,7 +62,7 @@ void PadTilingFunc(const ge::Shape srcShape, const ge::Shape oriSrcShape, const
62	END_TILING_DATA_DEF;	62	END_TILING_DATA_DEF;
63	```	63	```
64		64
65	-2. Tiling实现函数中，首先调用GetPadMaxMinTmpSize接口获取Pad接口能完成计算所需最大/最小临时空间大小，根据该范围结合实际的内存使用情况设置合适的空间大小；然后根据输入shape、剩余的可供计算的空间大小等信息获取Pad kernel侧接口所需tiling参数。	65	+2. Tiling实现函数中，首先调用GetPadMaxMinTmpSize接口获取Pad接口能完成计算所需最大和最小临时空间大小，根据该范围结合实际的内存使用情况设置合适的空间大小；然后根据输入shape、剩余的可供计算的空间大小等信息获取Pad kernel侧接口所需tiling参数。
66		66
67	```	67	```
68	namespace optiling {	68	namespace optiling {
	@@ -85,7 +85,7 @@ void PadTilingFunc(const ge::Shape srcShape, const ge::Shape oriSrcShape, const
85	uint32_t maxValue = 0;	85	uint32_t maxValue = 0;
86	uint32_t minValue = 0;	86	uint32_t minValue = 0;
87	AscendC::GetPadMaxMinTmpSize(srcShape, sizeof(half), maxValue, minValue);	87	AscendC::GetPadMaxMinTmpSize(srcShape, sizeof(half), maxValue, minValue);
88	- // 本样例中仅作为样例说明，获取最小值并传入，来保证功能正确，开发者可以根据需要传入合适的空间大小	88	+ // 本样例仅用于说明，获取最小值并传入，来保证功能正确，开发者可以根据需要传入合适的空间大小
89	const uint32_t localWorkSpaceSize = minValue;	89	const uint32_t localWorkSpaceSize = minValue;
90	AscendC::PadTilingFunc(srcShape, oriSrcShape, localWorkSpaceSize , sizeof(half), tiling.padTilingData);	90	AscendC::PadTilingFunc(srcShape, oriSrcShape, localWorkSpaceSize , sizeof(half), tiling.padTilingData);
91	// 其他逻辑	91	// 其他逻辑

Mdocs/api/SIMD-API/高阶API/张量变换/Transpose-Tiling.md+7-5

	@@ -89,7 +89,7 @@
89	END_TILING_DATA_DEF;	89	END_TILING_DATA_DEF;
90	```	90	```
91		91
92	-2. Tiling实现函数中，根据输入shape、可供计算的空间大小\(stackBufferSize\)等信息获取Transpose kernel侧接口所需tiling参数。	92	+2. Tiling实现函数中，首先调用GetTransposeMaxMinTmpSize接口获取Transpose接口能完成计算所需最大和最小临时空间大小，根据该范围结合实际的内存使用情况设置合适的空间大小；然后根据输入shape、剩余的可供计算的空间大小等信息获取Transpose kernel侧接口所需Tiling参数。
93		93
94	```	94	```
95	namespace optiling {	95	namespace optiling {
	@@ -110,11 +110,13 @@
110	uint32_t maxValue = 0;	110	uint32_t maxValue = 0;
111	uint32_t minValue = 0;	111	uint32_t minValue = 0;
112	AscendC::GetTransposeMaxMinTmpSize(srcShape, sizeof(half), transposeTypeIn, maxValue, minValue);	112	AscendC::GetTransposeMaxMinTmpSize(srcShape, sizeof(half), transposeTypeIn, maxValue, minValue);
113	- // 本样例中仅作为样例说明，获取最小值并传入，来保证功能正确，开发者可以根据需要传入合适的空间大小	113	+ // 本样例为保证功能正确，传入最小临时空间大小，开发者可以根据需要传入合适的空间大小
114	const uint32_t stackBufferSize = minValue;	114	const uint32_t stackBufferSize = minValue;
115	- // 获取Transpose Tiling参数	115	+ // 获取Transpose接口所需的Tiling参数
116	- AscendC::GetTransposeTilingInfo(srcShape, stackBufferSize, sizeof(half), transposeTypeIn, tiling.confusionTransposeTilingData);	116	+ AscendC::GetTransposeTilingInfo(
117	- ... // 其他逻辑	117	+ srcShape, stackBufferSize, sizeof(half), transposeTypeIn, tiling.confusionTransposeTilingData);
		118	+ ...
		119	+ // 其他逻辑
118	tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());	120	tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
119	context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());	121	context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
120	context->SetTilingKey(1);	122	context->SetTilingKey(1);

Mdocs/api/SIMD-API/高阶API/张量变换/UnPad-Tiling.md+6-6

Mdocs/api/SIMD-API/高阶API/张量变换/UnPad.md+1-1

	@@ -78,7 +78,7 @@ struct UnPadParams {
78	```	78	```
79	// dstLocal：输出Tensor	79	// dstLocal：输出Tensor
80	// srcLocal：输入Tensor	80	// srcLocal：输入Tensor
81	-// unPadParams：控制填充的参数	81	+// unPadParams：控制去除填充的参数
82	AscendC::UnPadParams unPadParams{0, 1}; // 左边去掉0列，右边去掉1列，当前暂不支持左边unpad	82	AscendC::UnPadParams unPadParams{0, 1}; // 左边去掉0列，右边去掉1列，当前暂不支持左边unpad
83	AscendC::UnPad(dstLocal, srcLocal, unPadParams, tilingData.unpadTilingData);	83	AscendC::UnPad(dstLocal, srcLocal, unPadParams, tilingData.unpadTilingData);
84	```	84	```

Mimpl/basic_api/kernel_operator_gemm_base_impl.h+119-119

	@@ -62,12 +62,12 @@ __aicore__ inline bool CheckRange(std::pair<uint32_t, uint32_t>& range, const ui
62		62
63	template <typename T, typename U, typename S>	63	template <typename T, typename U, typename S>
64	__aicore__ inline bool CheckOverflow(const LocalTensor<T>& dst, const LocalTensor<U>& src0,	64	__aicore__ inline bool CheckOverflow(const LocalTensor<T>& dst, const LocalTensor<U>& src0,
65	- const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling& tilling)	65	+ const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling& tiling)
66	{	66	{
67	// check l0c	67	// check l0c
68	- uint32_t roundM = DivCeil(m, tilling.blockSize) * tilling.blockSize;	68	+ uint32_t roundM = DivCeil(m, tiling.blockSize) * tiling.blockSize;
69	- uint32_t roundN = DivCeil(n, tilling.blockSize) * tilling.blockSize;	69	+ uint32_t roundN = DivCeil(n, tiling.blockSize) * tiling.blockSize;
70	- uint32_t roundK = DivCeil(k, tilling.c0Size) * tilling.c0Size;	70	+ uint32_t roundK = DivCeil(k, tiling.c0Size) * tiling.c0Size;
71		71
72	uint32_t needElementLoc = roundM * roundN * sizeof(uint32_t);	72	uint32_t needElementLoc = roundM * roundN * sizeof(uint32_t);
73	if (needElementLoc > TOTAL_L0C_SIZE) {	73	if (needElementLoc > TOTAL_L0C_SIZE) {
	@@ -88,10 +88,10 @@ __aicore__ inline bool CheckOverflow(const LocalTensor<T>& dst, const LocalTenso
88		88
89	template <typename T, typename U, typename S>	89	template <typename T, typename U, typename S>
90	__aicore__ inline bool CheckParams(const LocalTensor<T>& dst, const LocalTensor<U>& src0,	90	__aicore__ inline bool CheckParams(const LocalTensor<T>& dst, const LocalTensor<U>& src0,
91	- const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling& tilling)	91	+ const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling& tiling)
92	{	92	{
93	// check c0Size	93	// check c0Size
94	- if (tilling.c0Size != 16 && tilling.c0Size != 32) {	94	+ if (tiling.c0Size != 16 && tiling.c0Size != 32) {
95	return false;	95	return false;
96	}	96	}
97	// check scope	97	// check scope
	@@ -126,7 +126,7 @@ __aicore__ inline bool CheckParams(const LocalTensor<T>& dst, const LocalTensor<
126	}	126	}
127		127
128	// check overflow	128	// check overflow
129	- if (!CheckOverflow(dst, src0, src1, m, k, n, tilling)) {	129	+ if (!CheckOverflow(dst, src0, src1, m, k, n, tiling)) {
130	return false;	130	return false;
131	}	131	}
132		132
	@@ -134,38 +134,38 @@ __aicore__ inline bool CheckParams(const LocalTensor<T>& dst, const LocalTensor<
134	}	134	}
135	#endif	135	#endif
136		136
137	-__aicore__ inline void CalculateGemmTiling(GemmTiling& tilling)	137	+__aicore__ inline void CalculateGemmTiling(GemmTiling& tiling)
138	{	138	{
139	- tilling.mIterNum = 1;	139	+ tiling.mIterNum = 1;
140	- tilling.nIterNum = 1;	140	+ tiling.nIterNum = 1;
141	- tilling.kIterNum = DivCeil(tilling.kBlockNum, tilling.kTileBlock);	141	+ tiling.kIterNum = DivCeil(tiling.kBlockNum, tiling.kTileBlock);
142		142
143	- tilling.mTileBlock = DivCeil(tilling.mBlockNum, tilling.mIterNum);	143	+ tiling.mTileBlock = DivCeil(tiling.mBlockNum, tiling.mIterNum);
144	- tilling.nTileBlock = DivCeil(tilling.nBlockNum, tilling.nIterNum);	144	+ tiling.nTileBlock = DivCeil(tiling.nBlockNum, tiling.nIterNum);
145		145
146	- tilling.kTailBlock = tilling.kBlockNum - (tilling.kIterNum - 1) * tilling.kTileBlock;	146	+ tiling.kTailBlock = tiling.kBlockNum - (tiling.kIterNum - 1) * tiling.kTileBlock;
147	- tilling.mTailBlock = tilling.mBlockNum - (tilling.mIterNum - 1) * tilling.mTileBlock; // mTailBlock <= mBlockNum	147	+ tiling.mTailBlock = tiling.mBlockNum - (tiling.mIterNum - 1) * tiling.mTileBlock; // mTailBlock <= mBlockNum
148	- tilling.nTailBlock = tilling.nBlockNum - (tilling.nIterNum - 1) * tilling.nTileBlock;	148	+ tiling.nTailBlock = tiling.nBlockNum - (tiling.nIterNum - 1) * tiling.nTileBlock;
149		149
150	- tilling.kHasTail = tilling.kTailBlock != tilling.kTileBlock;	150	+ tiling.kHasTail = tiling.kTailBlock != tiling.kTileBlock;
151	- tilling.kHasTailEle = tilling.roundK != tilling.kNum;	151	+ tiling.kHasTailEle = tiling.roundK != tiling.kNum;
152	- tilling.kTailEle = tilling.kNum % (tilling.kTileBlock * tilling.c0Size);	152	+ tiling.kTailEle = tiling.kNum % (tiling.kTileBlock * tiling.c0Size);
153		153
154	- if (tilling.mNum != tilling.mTileBlock * tilling.blockSize) {	154	+ if (tiling.mNum != tiling.mTileBlock * tiling.blockSize) {
155	- tilling.mHasTail = true;	155	+ tiling.mHasTail = true;
156	} else {	156	} else {
157	- tilling.mHasTail = false;	157	+ tiling.mHasTail = false;
158	}	158	}
159	- tilling.nHasTail = tilling.nTileBlock != tilling.nTailBlock;	159	+ tiling.nHasTail = tiling.nTileBlock != tiling.nTailBlock;
160	}	160	}
161		161
162	template <typename T>	162	template <typename T>
163	-__aicore__ inline void LoadL0B(uint32_t kBlocks, uint32_t nBlocks, GemmTiling tilling, uint32_t i, uint32_t j,	163	+__aicore__ inline void LoadL0B(uint32_t kBlocks, uint32_t nBlocks, GemmTiling tiling, uint32_t i, uint32_t j,
164	const LocalTensor<T>& src1, const LocalTensor<T>& l0b)	164	const LocalTensor<T>& src1, const LocalTensor<T>& l0b)
165	{	165	{
166	- if (tilling.nIterNum == 1) {	166	+ if (tiling.nIterNum == 1) {
167	- uint32_t wSize = tilling.blockSize * tilling.c0Size;	167	+ uint32_t wSize = tiling.blockSize * tiling.c0Size;
168	- uint32_t wIdx = (i * tilling.kTileBlock * tilling.nBlockNum + j * tilling.nTileBlock) * wSize;	168	+ uint32_t wIdx = (i * tiling.kTileBlock * tiling.nBlockNum + j * tiling.nTileBlock) * wSize;
169	LoadData2DParams params;	169	LoadData2DParams params;
170	params.startIndex = 0;	170	params.startIndex = 0;
171	params.repeatTimes = kBlocks * nBlocks;	171	params.repeatTimes = kBlocks * nBlocks;
	@@ -174,10 +174,10 @@ __aicore__ inline void LoadL0B(uint32_t kBlocks, uint32_t nBlocks, GemmTiling ti
174	} else {	174	} else {
175	// load data row by row	175	// load data row by row
176	for (size_t index = 0; index < kBlocks; ++index) {	176	for (size_t index = 0; index < kBlocks; ++index) {
177	- uint32_t wSize = j * tilling.nTileBlock * tilling.blockSize * tilling.c0Size;	177	+ uint32_t wSize = j * tiling.nTileBlock * tiling.blockSize * tiling.c0Size;
178	uint32_t wIdx =	178	uint32_t wIdx =
179	- (i * tilling.kTileBlock + index) * tilling.nBlockNum * tilling.blockSize * tilling.c0Size + wSize;	179	+ (i * tiling.kTileBlock + index) * tiling.nBlockNum * tiling.blockSize * tiling.c0Size + wSize;
180	- uint32_t l0bIdx = index * nBlocks * tilling.blockSize * tilling.c0Size;	180	+ uint32_t l0bIdx = index * nBlocks * tiling.blockSize * tiling.c0Size;
181	LoadData2DParams params;	181	LoadData2DParams params;
182	params.startIndex = 0;	182	params.startIndex = 0;
183	params.repeatTimes = nBlocks;	183	params.repeatTimes = nBlocks;
	@@ -188,12 +188,12 @@ __aicore__ inline void LoadL0B(uint32_t kBlocks, uint32_t nBlocks, GemmTiling ti
188	}	188	}
189		189
190	template <typename T>	190	template <typename T>
191	-__aicore__ inline void LoadL0A(uint32_t kBlocks, uint32_t mBlocks, GemmTiling tilling, uint32_t i, uint32_t t,	191	+__aicore__ inline void LoadL0A(uint32_t kBlocks, uint32_t mBlocks, GemmTiling tiling, uint32_t i, uint32_t t,
192	const LocalTensor<T>& src0, const LocalTensor<T>& l0a)	192	const LocalTensor<T>& src0, const LocalTensor<T>& l0a)
193	{	193	{
194	if (kBlocks == 1) {	194	if (kBlocks == 1) {
195	- uint32_t l1aSize = i * tilling.kTileBlock * tilling.mBlockNum * tilling.blockSize * tilling.c0Size;	195	+ uint32_t l1aSize = i * tiling.kTileBlock * tiling.mBlockNum * tiling.blockSize * tiling.c0Size;
196	- uint32_t l1aOffset = t * tilling.mTileBlock * tilling.blockSize * tilling.c0Size + l1aSize;	196	+ uint32_t l1aOffset = t * tiling.mTileBlock * tiling.blockSize * tiling.c0Size + l1aSize;
197	LoadData2DParams params;	197	LoadData2DParams params;
198	params.startIndex = 0;	198	params.startIndex = 0;
199	params.repeatTimes = mBlocks;	199	params.repeatTimes = mBlocks;
	@@ -202,13 +202,13 @@ __aicore__ inline void LoadL0A(uint32_t kBlocks, uint32_t mBlocks, GemmTiling ti
202	} else {	202	} else {
203	// load data row by row	203	// load data row by row
204	for (size_t index = 0; index < mBlocks; index++) {	204	for (size_t index = 0; index < mBlocks; index++) {
205	- uint32_t l0aOffset = index * kBlocks * tilling.blockSize * tilling.c0Size;	205	+ uint32_t l0aOffset = index * kBlocks * tiling.blockSize * tiling.c0Size;
206	- uint32_t l1aOffset = (t * tilling.mTileBlock + index) * tilling.blockSize * tilling.c0Size +	206	+ uint32_t l1aOffset = (t * tiling.mTileBlock + index) * tiling.blockSize * tiling.c0Size +
207	- i * tilling.kTileBlock * tilling.mBlockNum * tilling.blockSize * tilling.c0Size;	207	+ i * tiling.kTileBlock * tiling.mBlockNum * tiling.blockSize * tiling.c0Size;
208	LoadData2DParams params;	208	LoadData2DParams params;
209	params.startIndex = 0;	209	params.startIndex = 0;
210	params.repeatTimes = kBlocks;	210	params.repeatTimes = kBlocks;
211	- params.srcStride = tilling.mBlockNum;	211	+ params.srcStride = tiling.mBlockNum;
212	LoadDataImpl(l0a[l0aOffset], src0[l1aOffset], params);	212	LoadDataImpl(l0a[l0aOffset], src0[l1aOffset], params);
213	}	213	}
214	}	214	}
	@@ -216,45 +216,45 @@ __aicore__ inline void LoadL0A(uint32_t kBlocks, uint32_t mBlocks, GemmTiling ti
216		216
217	template <typename T, typename U, typename S>	217	template <typename T, typename U, typename S>
218	__aicore__ inline void MmadFunc(const LocalTensor<U>& l0a, const LocalTensor<S>& l0b,	218	__aicore__ inline void MmadFunc(const LocalTensor<U>& l0a, const LocalTensor<S>& l0b,
219	- const LocalTensor<T>& l0c, int32_t initValue, GemmTiling tilling, size_t i)	219	+ const LocalTensor<T>& l0c, int32_t initValue, GemmTiling tiling, size_t i)
220	{	220	{
221	MmadParams mmadParams;	221	MmadParams mmadParams;
222	- mmadParams.m = tilling.mTileBlock * tilling.blockSize;	222	+ mmadParams.m = tiling.mTileBlock * tiling.blockSize;
223	- mmadParams.n = tilling.nTileBlock * tilling.blockSize;	223	+ mmadParams.n = tiling.nTileBlock * tiling.blockSize;
224	mmadParams.isBias = 1;	224	mmadParams.isBias = 1;
225		225
226	- if (tilling.kIterNum == 1) {	226	+ if (tiling.kIterNum == 1) {
227	- mmadParams.k = tilling.kNum;	227	+ mmadParams.k = tiling.kNum;
228	mmadParams.isBias = initValue;	228	mmadParams.isBias = initValue;
229	- } else if (initValue == 1 && tilling.kHasTailEle) {	229	+ } else if (initValue == 1 && tiling.kHasTailEle) {
230	- if (i == tilling.kIterNum - 1) {	230	+ if (i == tiling.kIterNum - 1) {
231	- mmadParams.k = tilling.kTailEle;	231	+ mmadParams.k = tiling.kTailEle;
232	} else {	232	} else {
233	- mmadParams.k = tilling.kTileBlock * tilling.c0Size;	233	+ mmadParams.k = tiling.kTileBlock * tiling.c0Size;
234	}	234	}
235	- } else if (initValue != 1 && tilling.kHasTailEle) {	235	+ } else if (initValue != 1 && tiling.kHasTailEle) {
236	if (i == 0) {	236	if (i == 0) {
237	- mmadParams.k = tilling.kTileBlock * tilling.c0Size;	237	+ mmadParams.k = tiling.kTileBlock * tiling.c0Size;
238	mmadParams.isBias = 0;	238	mmadParams.isBias = 0;
239	- } else if (i == tilling.kIterNum - 1) {	239	+ } else if (i == tiling.kIterNum - 1) {
240	- mmadParams.k = tilling.kTailEle;	240	+ mmadParams.k = tiling.kTailEle;
241	} else {	241	} else {
242	- mmadParams.k = tilling.kTileBlock * tilling.c0Size;	242	+ mmadParams.k = tiling.kTileBlock * tiling.c0Size;
243	}	243	}
244	- } else if (initValue == 1 && !tilling.kHasTailEle) {	244	+ } else if (initValue == 1 && !tiling.kHasTailEle) {
245	- if (i == tilling.kIterNum - 1) {	245	+ if (i == tiling.kIterNum - 1) {
246	- mmadParams.k = tilling.kTailBlock * tilling.c0Size;	246	+ mmadParams.k = tiling.kTailBlock * tiling.c0Size;
247	} else {	247	} else {
248	- mmadParams.k = tilling.kTileBlock * tilling.c0Size;	248	+ mmadParams.k = tiling.kTileBlock * tiling.c0Size;
249	}	249	}
250	} else {	250	} else {
251	if (i == 0) {	251	if (i == 0) {
252	- mmadParams.k = tilling.kTileBlock * tilling.c0Size;	252	+ mmadParams.k = tiling.kTileBlock * tiling.c0Size;
253	mmadParams.isBias = 0;	253	mmadParams.isBias = 0;
254	- } else if (i == tilling.kIterNum - 1) {	254	+ } else if (i == tiling.kIterNum - 1) {
255	- mmadParams.k = tilling.kTailBlock * tilling.c0Size;	255	+ mmadParams.k = tiling.kTailBlock * tiling.c0Size;
256	} else {	256	} else {
257	- mmadParams.k = tilling.kTileBlock * tilling.c0Size;	257	+ mmadParams.k = tiling.kTileBlock * tiling.c0Size;
258	}	258	}
259	}	259	}
260	MmadImpl(l0c, l0a, l0b, mmadParams);	260	MmadImpl(l0c, l0a, l0b, mmadParams);
	@@ -307,30 +307,30 @@ __aicore__ inline void GetSingleThreadBuffer(LocalTensor<T>& l0a, LocalTensor<U>
307		307
308	template <typename T, typename U, typename S>	308	template <typename T, typename U, typename S>
309	__aicore__ inline void GemmExecNmNopingpong(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,	309	__aicore__ inline void GemmExecNmNopingpong(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,
310	- const LocalTensor<S>& src1, GemmTiling tilling, const int32_t initValue)	310	+ const LocalTensor<S>& src1, GemmTiling tiling, const int32_t initValue)
311	{	311	{
312	LocalTensor<U> l0a;	312	LocalTensor<U> l0a;
313	LocalTensor<S> l0b;	313	LocalTensor<S> l0b;
314	GetSingleThreadBuffer(l0a, l0b);	314	GetSingleThreadBuffer(l0a, l0b);
315	event_t eventIdMToMte1 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::M_MTE1));	315	event_t eventIdMToMte1 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::M_MTE1));
316	SetFlag<HardEvent::M_MTE1>(eventIdMToMte1);	316	SetFlag<HardEvent::M_MTE1>(eventIdMToMte1);
317	- for (size_t indexK = 0; indexK < tilling.kIterNum; indexK++) {	317	+ for (size_t indexK = 0; indexK < tiling.kIterNum; indexK++) {
318	- uint32_t kBlocks = tilling.kTileBlock;	318	+ uint32_t kBlocks = tiling.kTileBlock;
319	- if (indexK == tilling.kIterNum - 1) {	319	+ if (indexK == tiling.kIterNum - 1) {
320	- kBlocks = tilling.kTailBlock;	320	+ kBlocks = tiling.kTailBlock;
321	}	321	}
322	WaitFlag<HardEvent::M_MTE1>(eventIdMToMte1);	322	WaitFlag<HardEvent::M_MTE1>(eventIdMToMte1);
323	- for (size_t indexN = 0; indexN < tilling.nIterNum; indexN++) {	323	+ for (size_t indexN = 0; indexN < tiling.nIterNum; indexN++) {
324	// load data from l1 to l0b	324	// load data from l1 to l0b
325	- LoadL0B(kBlocks, tilling.nTileBlock, tilling, indexK, indexN, src1, l0b);	325	+ LoadL0B(kBlocks, tiling.nTileBlock, tiling, indexK, indexN, src1, l0b);
326	- for (size_t indexM = 0; indexM < tilling.mIterNum; indexM++) {	326	+ for (size_t indexM = 0; indexM < tiling.mIterNum; indexM++) {
327	// load data from l1 to l0a	327	// load data from l1 to l0a
328	- LoadL0A(kBlocks, tilling.mTileBlock, tilling, indexK, indexM, src0, l0a);	328	+ LoadL0A(kBlocks, tiling.mTileBlock, tiling, indexK, indexM, src0, l0a);
329	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));	329	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));
330	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	330	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
331	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	331	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
332	PipeBarrier<PIPE_M>();	332	PipeBarrier<PIPE_M>();
333	- MmadFunc(l0a, l0b, l0c, initValue, tilling, indexK);	333	+ MmadFunc(l0a, l0b, l0c, initValue, tiling, indexK);
334	}	334	}
335	}	335	}
336	SetFlag<HardEvent::M_MTE1>(eventIdMToMte1);	336	SetFlag<HardEvent::M_MTE1>(eventIdMToMte1);
	@@ -340,7 +340,7 @@ __aicore__ inline void GemmExecNmNopingpong(const LocalTensor<T>& l0c, const Loc
340		340
341	template <typename T, typename U, typename S>	341	template <typename T, typename U, typename S>
342	__aicore__ inline void GemmExecNmPingPong(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,	342	__aicore__ inline void GemmExecNmPingPong(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,
343	- const LocalTensor<S>& src1, GemmTiling tilling, const int32_t initValue)	343	+ const LocalTensor<S>& src1, GemmTiling tiling, const int32_t initValue)
344	{	344	{
345	uint32_t ping = 1;	345	uint32_t ping = 1;
346	LocalTensor<U> l0aPing;	346	LocalTensor<U> l0aPing;
	@@ -354,40 +354,40 @@ __aicore__ inline void GemmExecNmPingPong(const LocalTensor<T>& l0c, const Local
354	SetFlag<HardEvent::M_MTE1>(eventId0);	354	SetFlag<HardEvent::M_MTE1>(eventId0);
355	SetFlag<HardEvent::M_MTE1>(eventId1);	355	SetFlag<HardEvent::M_MTE1>(eventId1);
356		356
357	- for (size_t i = 0; i < tilling.kIterNum; i++) {	357	+ for (size_t i = 0; i < tiling.kIterNum; i++) {
358	- uint32_t kBlocks = tilling.kTileBlock;	358	+ uint32_t kBlocks = tiling.kTileBlock;
359	- if (i == tilling.kIterNum - 1) {	359	+ if (i == tiling.kIterNum - 1) {
360	- kBlocks = tilling.kTailBlock;	360	+ kBlocks = tiling.kTailBlock;
361	}	361	}
362	if (ping == 1) {	362	if (ping == 1) {
363	WaitFlag<HardEvent::M_MTE1>(eventId0);	363	WaitFlag<HardEvent::M_MTE1>(eventId0);
364	- for (size_t indexN = 0; indexN < tilling.nIterNum; indexN++) {	364	+ for (size_t indexN = 0; indexN < tiling.nIterNum; indexN++) {
365	// load data from l1 to l0b	365	// load data from l1 to l0b
366	- LoadL0B(kBlocks, tilling.nTileBlock, tilling, i, indexN, src1, l0bPing);	366	+ LoadL0B(kBlocks, tiling.nTileBlock, tiling, i, indexN, src1, l0bPing);
367	- for (size_t indexM = 0; indexM < tilling.mIterNum; indexM++) {	367	+ for (size_t indexM = 0; indexM < tiling.mIterNum; indexM++) {
368	// load data from l1 to l0a	368	// load data from l1 to l0a
369	- LoadL0A(kBlocks, tilling.mTileBlock, tilling, i, indexM, src0, l0aPing);	369	+ LoadL0A(kBlocks, tiling.mTileBlock, tiling, i, indexM, src0, l0aPing);
370	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));	370	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));
371	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	371	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
372	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	372	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
373	PipeBarrier<PIPE_M>();	373	PipeBarrier<PIPE_M>();
374	- MmadFunc(l0aPing, l0bPing, l0c, initValue, tilling, i);	374	+ MmadFunc(l0aPing, l0bPing, l0c, initValue, tiling, i);
375	}	375	}
376	}	376	}
377	SetFlag<HardEvent::M_MTE1>(eventId0);	377	SetFlag<HardEvent::M_MTE1>(eventId0);
378	} else {	378	} else {
379	WaitFlag<HardEvent::M_MTE1>(eventId1);	379	WaitFlag<HardEvent::M_MTE1>(eventId1);
380	- for (size_t indexN = 0; indexN < tilling.nIterNum; indexN++) {	380	+ for (size_t indexN = 0; indexN < tiling.nIterNum; indexN++) {
381	// load data from l1 to l0b	381	// load data from l1 to l0b
382	- LoadL0B(kBlocks, tilling.nTileBlock, tilling, i, indexN, src1, l0bPong);	382	+ LoadL0B(kBlocks, tiling.nTileBlock, tiling, i, indexN, src1, l0bPong);
383	- for (size_t indexM = 0; indexM < tilling.mIterNum; indexM++) {	383	+ for (size_t indexM = 0; indexM < tiling.mIterNum; indexM++) {
384	// load data from l1 to l0a	384	// load data from l1 to l0a
385	- LoadL0A(kBlocks, tilling.mTileBlock, tilling, i, indexM, src0, l0aPong);	385	+ LoadL0A(kBlocks, tiling.mTileBlock, tiling, i, indexM, src0, l0aPong);
386	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));	386	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));
387	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	387	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
388	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	388	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
389	PipeBarrier<PIPE_M>();	389	PipeBarrier<PIPE_M>();
390	- MmadFunc(l0aPong, l0bPong, l0c, initValue, tilling, i);	390	+ MmadFunc(l0aPong, l0bPong, l0c, initValue, tiling, i);
391	}	391	}
392	}	392	}
393	SetFlag<HardEvent::M_MTE1>(eventId1);	393	SetFlag<HardEvent::M_MTE1>(eventId1);
	@@ -408,43 +408,43 @@ __aicore__ inline void GemmExecNmPingPong(const LocalTensor<T>& l0c, const Local
408		408
409	template <typename T, typename U, typename S>	409	template <typename T, typename U, typename S>
410	__aicore__ inline void GemmExecNm(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,	410	__aicore__ inline void GemmExecNm(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,
411	- const LocalTensor<S>& src1, GemmTiling tilling, const int32_t initValue)	411	+ const LocalTensor<S>& src1, GemmTiling tiling, const int32_t initValue)
412	{	412	{
413	- uint32_t needL0Asize = tilling.roundM * tilling.dtypeSize * tilling.c0Size * tilling.kTileBlock * 2;	413	+ uint32_t needL0Asize = tiling.roundM * tiling.dtypeSize * tiling.c0Size * tiling.kTileBlock * 2;
414	- uint32_t needL0Bsize = tilling.roundN * tilling.dtypeSize * tilling.c0Size * tilling.kTileBlock * 2;	414	+ uint32_t needL0Bsize = tiling.roundN * tiling.dtypeSize * tiling.c0Size * tiling.kTileBlock * 2;
415	if (needL0Asize > TOTAL_L0A_SIZE \|\| needL0Bsize > TOTAL_L0B_SIZE) {	415	if (needL0Asize > TOTAL_L0A_SIZE \|\| needL0Bsize > TOTAL_L0B_SIZE) {
416	- GemmExecNmNopingpong(l0c, src0, src1, tilling, initValue);	416	+ GemmExecNmNopingpong(l0c, src0, src1, tiling, initValue);
417	return;	417	return;
418	}	418	}
419	- GemmExecNmPingPong(l0c, src0, src1, tilling, initValue);	419	+ GemmExecNmPingPong(l0c, src0, src1, tiling, initValue);
420	}	420	}
421		421
422	template <typename T, typename U, typename S>	422	template <typename T, typename U, typename S>
423	__aicore__ inline void GemmExecMnNopingpong(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,	423	__aicore__ inline void GemmExecMnNopingpong(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,
424	- const LocalTensor<S>& src1, GemmTiling tilling, const int32_t initValue)	424	+ const LocalTensor<S>& src1, GemmTiling tiling, const int32_t initValue)
425	{	425	{
426	LocalTensor<S> l0b;	426	LocalTensor<S> l0b;
427	LocalTensor<U> l0a;	427	LocalTensor<U> l0a;
428	GetSingleThreadBuffer(l0a, l0b);	428	GetSingleThreadBuffer(l0a, l0b);
429	event_t eventIdMToMte1 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::M_MTE1));	429	event_t eventIdMToMte1 = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::M_MTE1));
430	SetFlag<HardEvent::M_MTE1>(eventIdMToMte1);	430	SetFlag<HardEvent::M_MTE1>(eventIdMToMte1);
431	- for (size_t indexK = 0; indexK < tilling.kIterNum; indexK++) {	431	+ for (size_t indexK = 0; indexK < tiling.kIterNum; indexK++) {
432	- uint32_t kBlocks = tilling.kTileBlock;	432	+ uint32_t kBlocks = tiling.kTileBlock;
433	- if (indexK == tilling.kIterNum - 1) {	433	+ if (indexK == tiling.kIterNum - 1) {
434	- kBlocks = tilling.kTailBlock;	434	+ kBlocks = tiling.kTailBlock;
435	}	435	}
436	WaitFlag<HardEvent::M_MTE1>(eventIdMToMte1);	436	WaitFlag<HardEvent::M_MTE1>(eventIdMToMte1);
437	- for (size_t indexM = 0; indexM < tilling.mIterNum; indexM++) {	437	+ for (size_t indexM = 0; indexM < tiling.mIterNum; indexM++) {
438	// load data from l1 to l0a	438	// load data from l1 to l0a
439	- LoadL0A(kBlocks, tilling.mTileBlock, tilling, indexK, indexM, src0, l0a);	439	+ LoadL0A(kBlocks, tiling.mTileBlock, tiling, indexK, indexM, src0, l0a);
440	- for (size_t indexN = 0; indexN < tilling.nIterNum; indexN++) {	440	+ for (size_t indexN = 0; indexN < tiling.nIterNum; indexN++) {
441	// load data from l1 to l0b	441	// load data from l1 to l0b
442	- LoadL0B(kBlocks, tilling.nTileBlock, tilling, indexK, indexN, src1, l0b);	442	+ LoadL0B(kBlocks, tiling.nTileBlock, tiling, indexK, indexN, src1, l0b);
443	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));	443	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));
444	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	444	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
445	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	445	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
446	PipeBarrier<PIPE_M>();	446	PipeBarrier<PIPE_M>();
447	- MmadFunc(l0a, l0b, l0c, initValue, tilling, indexK);	447	+ MmadFunc(l0a, l0b, l0c, initValue, tiling, indexK);
448	}	448	}
449	}	449	}
450	SetFlag<HardEvent::M_MTE1>(eventIdMToMte1);	450	SetFlag<HardEvent::M_MTE1>(eventIdMToMte1);
	@@ -454,7 +454,7 @@ __aicore__ inline void GemmExecMnNopingpong(const LocalTensor<T>& l0c, const Loc
454		454
455	template <typename T, typename U, typename S>	455	template <typename T, typename U, typename S>
456	__aicore__ inline void GemmExecMnPingPong(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,	456	__aicore__ inline void GemmExecMnPingPong(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,
457	- const LocalTensor<S>& src1, GemmTiling tilling, const int32_t initValue)	457	+ const LocalTensor<S>& src1, GemmTiling tiling, const int32_t initValue)
458	{	458	{
459	uint32_t ping = 1;	459	uint32_t ping = 1;
460	LocalTensor<U> l0aPing;	460	LocalTensor<U> l0aPing;
	@@ -468,41 +468,41 @@ __aicore__ inline void GemmExecMnPingPong(const LocalTensor<T>& l0c, const Local
468	SetFlag<HardEvent::M_MTE1>(eventId0);	468	SetFlag<HardEvent::M_MTE1>(eventId0);
469	SetFlag<HardEvent::M_MTE1>(eventId1);	469	SetFlag<HardEvent::M_MTE1>(eventId1);
470		470
471	- for (size_t i = 0; i < tilling.kIterNum; i++) {	471	+ for (size_t i = 0; i < tiling.kIterNum; i++) {
472	- uint32_t kBlocks = tilling.kTileBlock;	472	+ uint32_t kBlocks = tiling.kTileBlock;
473	- if (i == tilling.kIterNum - 1) {	473	+ if (i == tiling.kIterNum - 1) {
474	- kBlocks = tilling.kTailBlock;	474	+ kBlocks = tiling.kTailBlock;
475	}	475	}
476	if (ping == 1) {	476	if (ping == 1) {
477	WaitFlag<HardEvent::M_MTE1>(eventId0);	477	WaitFlag<HardEvent::M_MTE1>(eventId0);
478	- for (size_t indexM = 0; indexM < tilling.mIterNum; indexM++) {	478	+ for (size_t indexM = 0; indexM < tiling.mIterNum; indexM++) {
479	// load data from l1 to l0a	479	// load data from l1 to l0a
480	- LoadL0A(kBlocks, tilling.mTileBlock, tilling, i, indexM, src0, l0aPing);	480	+ LoadL0A(kBlocks, tiling.mTileBlock, tiling, i, indexM, src0, l0aPing);
481	- for (size_t indexN = 0; indexN < tilling.nIterNum; indexN++) {	481	+ for (size_t indexN = 0; indexN < tiling.nIterNum; indexN++) {
482	// load data from l1 to l0b	482	// load data from l1 to l0b
483	- LoadL0B(kBlocks, tilling.nTileBlock, tilling, i, indexN, src1, l0bPing);	483	+ LoadL0B(kBlocks, tiling.nTileBlock, tiling, i, indexN, src1, l0bPing);
484		484
485	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));	485	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));
486	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	486	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
487	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	487	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
488	PipeBarrier<PIPE_M>();	488	PipeBarrier<PIPE_M>();
489	- MmadFunc(l0aPing, l0bPing, l0c, initValue, tilling, i);	489	+ MmadFunc(l0aPing, l0bPing, l0c, initValue, tiling, i);
490	}	490	}
491	}	491	}
492	SetFlag<HardEvent::M_MTE1>(eventId0);	492	SetFlag<HardEvent::M_MTE1>(eventId0);
493	} else {	493	} else {
494	WaitFlag<HardEvent::M_MTE1>(eventId1);	494	WaitFlag<HardEvent::M_MTE1>(eventId1);
495	- for (size_t indexM = 0; indexM < tilling.mIterNum; indexM++) {	495	+ for (size_t indexM = 0; indexM < tiling.mIterNum; indexM++) {
496	// load data from l1 to l0a	496	// load data from l1 to l0a
497	- LoadL0A(kBlocks, tilling.mTileBlock, tilling, i, indexM, src0, l0aPong);	497	+ LoadL0A(kBlocks, tiling.mTileBlock, tiling, i, indexM, src0, l0aPong);
498	- for (size_t indexN = 0; indexN < tilling.nIterNum; indexN++) {	498	+ for (size_t indexN = 0; indexN < tiling.nIterNum; indexN++) {
499	// load data from l1 to l0b	499	// load data from l1 to l0b
500	- LoadL0B(kBlocks, tilling.nTileBlock, tilling, i, indexN, src1, l0bPong);	500	+ LoadL0B(kBlocks, tiling.nTileBlock, tiling, i, indexN, src1, l0bPong);
501	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));	501	event_t eventIdMte1ToM = static_cast<event_t>(GetTPipePtr()->FetchEventID(HardEvent::MTE1_M));
502	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	502	SetFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
503	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);	503	WaitFlag<HardEvent::MTE1_M>(eventIdMte1ToM);
504	PipeBarrier<PIPE_M>();	504	PipeBarrier<PIPE_M>();
505	- MmadFunc(l0aPong, l0bPong, l0c, initValue, tilling, i);	505	+ MmadFunc(l0aPong, l0bPong, l0c, initValue, tiling, i);
506	}	506	}
507	}	507	}
508	SetFlag<HardEvent::M_MTE1>(eventId1);	508	SetFlag<HardEvent::M_MTE1>(eventId1);
	@@ -518,15 +518,15 @@ __aicore__ inline void GemmExecMnPingPong(const LocalTensor<T>& l0c, const Local
518		518
519	template <typename T, typename U, typename S>	519	template <typename T, typename U, typename S>
520	__aicore__ inline void GemmExecMn(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,	520	__aicore__ inline void GemmExecMn(const LocalTensor<T>& l0c, const LocalTensor<U>& src0,
521	- const LocalTensor<S>& src1, GemmTiling tilling, const int32_t initValue)	521	+ const LocalTensor<S>& src1, GemmTiling tiling, const int32_t initValue)
522	{	522	{
523	- uint32_t needL0Bsize = tilling.roundN * tilling.dtypeSize * tilling.c0Size * tilling.kTileBlock * 2;	523	+ uint32_t needL0Bsize = tiling.roundN * tiling.dtypeSize * tiling.c0Size * tiling.kTileBlock * 2;
524	- uint32_t needL0Asize = tilling.roundM * tilling.dtypeSize * tilling.c0Size * tilling.kTileBlock * 2;	524	+ uint32_t needL0Asize = tiling.roundM * tiling.dtypeSize * tiling.c0Size * tiling.kTileBlock * 2;
525	if (needL0Asize > TOTAL_L0A_SIZE \|\| needL0Bsize > TOTAL_L0B_SIZE) {	525	if (needL0Asize > TOTAL_L0A_SIZE \|\| needL0Bsize > TOTAL_L0B_SIZE) {
526	- GemmExecMnNopingpong(l0c, src0, src1, tilling, initValue);	526	+ GemmExecMnNopingpong(l0c, src0, src1, tiling, initValue);
527	return;	527	return;
528	}	528	}
529	- GemmExecMnPingPong(l0c, src0, src1, tilling, initValue);	529	+ GemmExecMnPingPong(l0c, src0, src1, tiling, initValue);
530	}	530	}
531	} // namespace AscendC	531	} // namespace AscendC
532	#endif // ASCENDC_MODULE_OPERATOR_GEMM_BASE_IMPL_H	532	#endif // ASCENDC_MODULE_OPERATOR_GEMM_BASE_IMPL_H

Mimpl/basic_api/kernel_operator_gemm_intf_impl.h+41-41

	@@ -40,33 +40,33 @@ __aicore__ inline GemmTiling GetGemmTiling(uint32_t m, uint32_t k, uint32_t n)
40	c0 = 16;	40	c0 = 16;
41	dSize = 2;	41	dSize = 2;
42	}	42	}
43	- GemmTiling tilling;	43	+ GemmTiling tiling;
44	- tilling.c0Size = c0;	44	+ tiling.c0Size = c0;
45	- tilling.dtypeSize = dSize;	45	+ tiling.dtypeSize = dSize;
46	- tilling.mNum = m;	46	+ tiling.mNum = m;
47	- tilling.nNum = n;	47	+ tiling.nNum = n;
48	- tilling.kNum = k;	48	+ tiling.kNum = k;
49	- tilling.roundM = DivCeil(m, tilling.blockSize) * tilling.blockSize; // blockSize = 16 * 16	49	+ tiling.roundM = DivCeil(m, tiling.blockSize) * tiling.blockSize; // blockSize = 16 * 16
50	- tilling.roundN = DivCeil(n, tilling.blockSize) * tilling.blockSize;	50	+ tiling.roundN = DivCeil(n, tiling.blockSize) * tiling.blockSize;
51	- tilling.roundK = DivCeil(k, tilling.c0Size) * tilling.c0Size; // c0Size = 16 \|\| c0Size = 32	51	+ tiling.roundK = DivCeil(k, tiling.c0Size) * tiling.c0Size; // c0Size = 16 \|\| c0Size = 32
52	- uint32_t k0a = TOTAL_L0A_SIZE / 2 / (tilling.roundM * dSize);	52	+ uint32_t k0a = TOTAL_L0A_SIZE / 2 / (tiling.roundM * dSize);
53	- uint32_t k0b = TOTAL_L0B_SIZE / 2 / (tilling.roundN * dSize);	53	+ uint32_t k0b = TOTAL_L0B_SIZE / 2 / (tiling.roundN * dSize);
54	uint32_t k0 = k0a > k0b ? k0b : k0a;	54	uint32_t k0 = k0a > k0b ? k0b : k0a;
55	k0 = k0 > k ? k : k0;	55	k0 = k0 > k ? k : k0;
56		56
57	- tilling.kTileBlock = k0 / tilling.c0Size;	57	+ tiling.kTileBlock = k0 / tiling.c0Size;
58	- if (tilling.kTileBlock == 0) {	58	+ if (tiling.kTileBlock == 0) {
59	- tilling.kTileBlock = 1;	59	+ tiling.kTileBlock = 1;
60	}	60	}
61	- tilling.loopMode = LoopMode::MODE_NM;	61	+ tiling.loopMode = LoopMode::MODE_NM;
62		62
63	- tilling.mBlockNum = DivCeil(m, tilling.blockSize);	63	+ tiling.mBlockNum = DivCeil(m, tiling.blockSize);
64	- tilling.nBlockNum = DivCeil(n, tilling.blockSize);	64	+ tiling.nBlockNum = DivCeil(n, tiling.blockSize);
65	- tilling.kBlockNum = DivCeil(k, tilling.c0Size);	65	+ tiling.kBlockNum = DivCeil(k, tiling.c0Size);
66		66
67	- CalculateGemmTiling(tilling);	67	+ CalculateGemmTiling(tiling);
68		68
69	- return tilling;	69	+ return tiling;
70	}	70	}
71		71
72	/*	72	/*
	@@ -78,20 +78,20 @@ __aicore__ inline GemmTiling GetGemmTiling(uint32_t m, uint32_t k, uint32_t n)
78	* @param [in] m Number of rows of src0	78	* @param [in] m Number of rows of src0
79	* @param [in] n Number of rows of src1	79	* @param [in] n Number of rows of src1
80	* @param [in] k Number of columns of src1	80	* @param [in] k Number of columns of src1
81	- * @param [in] tilling.blockSize size of block	81	+ * @param [in] tiling.blockSize size of block
82	- * @param [in] tilling.mNum args of m	82	+ * @param [in] tiling.mNum args of m
83	- * @param [in] tilling.nNum args of n	83	+ * @param [in] tiling.nNum args of n
84	- * @param [in] tilling.kNum args of k	84	+ * @param [in] tiling.kNum args of k
85	- * @param [in] tilling.roundM/N/K Rounding parameter	85	+ * @param [in] tiling.roundM/N/K Rounding parameter
86	- * @param [in] tilling.c0Size The byte length of a block	86	+ * @param [in] tiling.c0Size The byte length of a block
87	- * @param [in] tilling.dtypeSize Byte length of the incoming data type	87	+ * @param [in] tiling.dtypeSize Byte length of the incoming data type
88	- * @param [in] tilling.m/n/kBlockNum Number of blocks of m/n/k axis	88	+ * @param [in] tiling.m/n/kBlockNum Number of blocks of m/n/k axis
89	- * @param [in] tilling.m/n/kIterNum Number of traversal dimensions	89	+ * @param [in] tiling.m/n/kIterNum Number of traversal dimensions
90	- * @param [in] tilling.m/k/nTileBlock Number of M/N/K axis cutting blocks	90	+ * @param [in] tiling.m/k/nTileBlock Number of M/N/K axis cutting blocks
91	- * @param [in] tilling.m/n/kHasTailNumber of tail blocks of M/K/N axis	91	+ * @param [in] tiling.m/n/kHasTailNumber of tail blocks of M/K/N axis
92	- * @param [in] tilling.kHasTileEle Judge whether the tail block exists	92	+ * @param [in] tiling.kHasTileEle Judge whether the tail block exists
93	- * @param [in] tilling.KtailEle K-axis tail block element	93	+ * @param [in] tiling.KtailEle K-axis tail block element
94	- * @param [in] tilling.kThreadNum K-axis passes	94	+ * @param [in] tiling.kThreadNum K-axis passes
95	* @param [in] partialsum judge whether the calculation result is moved out	95	* @param [in] partialsum judge whether the calculation result is moved out
96	* @param [in] initValue Initialization parameters	96	* @param [in] initValue Initialization parameters
97	*/	97	*/
	@@ -99,11 +99,11 @@ template <typename T, typename U, typename S>
99	[[deprecated("NOTICE: Gemm has been deprecated and will be removed in the next version. "	99	[[deprecated("NOTICE: Gemm has been deprecated and will be removed in the next version. "
100	"Please do not use it!")]]	100	"Please do not use it!")]]
101	__aicore__ inline __inout_pipe__(V) void Gemm(const LocalTensor<T>& dst, const LocalTensor<U>& src0,	101	__aicore__ inline __inout_pipe__(V) void Gemm(const LocalTensor<T>& dst, const LocalTensor<U>& src0,
102	- const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling tilling,	102	+ const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling tiling,
103	bool partialsum, int32_t initValue)	103	bool partialsum, int32_t initValue)
104	{	104	{
105	#if ASCENDC_CPU_DEBUG	105	#if ASCENDC_CPU_DEBUG
106	- bool flag = CheckParams(dst, src0, src1, m, k, n, tilling);	106	+ bool flag = CheckParams(dst, src0, src1, m, k, n, tiling);
107	if (!flag) {	107	if (!flag) {
108	return;	108	return;
109	}	109	}
	@@ -132,10 +132,10 @@ __aicore__ inline __inout_pipe__(V) void Gemm(const LocalTensor<T>& dst, const L
132	#endif	132	#endif
133	}	133	}
134		134
135	- if (tilling.loopMode == LoopMode::MODE_NM) {	135	+ if (tiling.loopMode == LoopMode::MODE_NM) {
136	- GemmExecNm(l0c, src0, src1, tilling, initValue);	136	+ GemmExecNm(l0c, src0, src1, tiling, initValue);
137	- } else if (tilling.loopMode == LoopMode::MODE_MN) {	137	+ } else if (tiling.loopMode == LoopMode::MODE_MN) {
138	- GemmExecMn(l0c, src0, src1, tilling, initValue);	138	+ GemmExecMn(l0c, src0, src1, tiling, initValue);
139	} else {	139	} else {
140	// other mode are not supported	140	// other mode are not supported
141	}	141	}
	@@ -143,7 +143,7 @@ __aicore__ inline __inout_pipe__(V) void Gemm(const LocalTensor<T>& dst, const L
143	#if (__NPU_ARCH__ == 1001) \|\| (__NPU_ARCH__ == 2002)	143	#if (__NPU_ARCH__ == 1001) \|\| (__NPU_ARCH__ == 2002)
144	if (dstScope == Hardware::UB) {	144	if (dstScope == Hardware::UB) {
145	pipe_barrier(PIPE_ALL);	145	pipe_barrier(PIPE_ALL);
146	- dataCopyParams.blockLen = tilling.roundM * tilling.roundN * sizeof(PrimT<T>) / 1024;	146	+ dataCopyParams.blockLen = tiling.roundM * tiling.roundN * sizeof(PrimT<T>) / 1024;
147	DataCopy(dst, l0c, dataCopyParams, enhancedParams);	147	DataCopy(dst, l0c, dataCopyParams, enhancedParams);
148	}	148	}
149	#endif	149	#endif

Minclude/basic_api/kernel_operator_gemm_intf.h+15-15

Mtests/api/basic_api/ascendc_case_common/test_operator_gemm.cpp+3-3

	@@ -59,9 +59,9 @@ extern "C" __global__ __aicore__ void main_gemm_test(__gm__ half* __restrict__ i
59	copy_gm_to_cbuf((__cbuf__ half)input0.GetPhyAddr(), input0Gm, 0, 1, roundm roundk * 2 / 32, 0, 0, PAD_NONE);	59	copy_gm_to_cbuf((__cbuf__ half)input0.GetPhyAddr(), input0Gm, 0, 1, roundm roundk * 2 / 32, 0, 0, PAD_NONE);
60	copy_gm_to_cbuf((__cbuf__ half)input1.GetPhyAddr(), input1Gm, 0, 1, roundk roundn * 2 / 32, 0, 0, PAD_NONE);	60	copy_gm_to_cbuf((__cbuf__ half)input1.GetPhyAddr(), input1Gm, 0, 1, roundk roundn * 2 / 32, 0, 0, PAD_NONE);
61	pipe_barrier(PIPE_ALL);	61	pipe_barrier(PIPE_ALL);
62	- GemmTiling tilling = GetGemmTiling<half>(m, k, n);	62	+ GemmTiling tiling = GetGemmTiling<half>(m, k, n);
63	- tilling.loopMode = mode;	63	+ tiling.loopMode = mode;
64	- Gemm(output, input0, input1, m, k, n, tilling, false, initValue);	64	+ Gemm(output, input0, input1, m, k, n, tiling, false, initValue);
65		65
66	pipe_barrier(PIPE_ALL);	66	pipe_barrier(PIPE_ALL);
67	copy_matrix_cc_to_ubuf((__ubuf__ float)ublocal.GetPhyAddr(), (__cc__ float)output.GetPhyAddr(), 0, 1,	67	copy_matrix_cc_to_ubuf((__ubuf__ float)ublocal.GetPhyAddr(), (__cc__ float)output.GetPhyAddr(), 0, 1,

Mtests/api/basic_api/ascendc_header_checker/kernel_operator_gemm_intf.cpp+8-5

	@@ -14,16 +14,19 @@
14	#include "kernel_operator_gemm_intf.h"	14	#include "kernel_operator_gemm_intf.h"
15	#endif	15	#endif
16		16
17	-// __aicore__ inline __inout_pipe__(V) void Gemm(const LocalTensor<T>& dst, const LocalTensor<U>& src0, const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling tilling, bool partialsum = true, int32_t initValue = 0);	17	+// __aicore__ inline __inout_pipe__(V) void Gemm(const LocalTensor<T>& dst, const LocalTensor<U>& src0, const
18	-extern "C" __global__ __aicore__ void KernelTestGemm1() {	18	+// LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling tiling, bool partialsum =
		19	+// true, int32_t initValue = 0);
		20	+extern "C" __global__ __aicore__ void KernelTestGemm1()
		21	+{
19	AscendC::LocalTensor<float> dst;	22	AscendC::LocalTensor<float> dst;
20	AscendC::LocalTensor<half> src0;	23	AscendC::LocalTensor<half> src0;
21	AscendC::LocalTensor<half> src1;	24	AscendC::LocalTensor<half> src1;
22	uint32_t m = 0;	25	uint32_t m = 0;
23	uint32_t k = 0;	26	uint32_t k = 0;
24	uint32_t n = 0;	27	uint32_t n = 0;
25	- AscendC::GemmTiling tilling;	28	+ AscendC::GemmTiling tiling;
26	- bool partialsum = true ;	29	+ bool partialsum = true;
27	int32_t initValue = 0;	30	int32_t initValue = 0;
28	- AscendC::Gemm(dst, src0, src1, m, k, n, tilling, true, 0);	31	+ AscendC::Gemm(dst, src0, src1, m, k, n, tiling, true, 0);
29	}	32	}

	@@ -164,7 +164,7 @@ Kirin X90，支持数据类型：int8_t、half。
164		164
165	如下图示例：	165	如下图示例：
166		166
167	- - 共需要处理3072字节的数据，每次迭代处理32\32\1字节数据，需要3次迭代可以完成，repeatTime = 3；	167	+ - 共需要处理3072字节的数据，每次迭代处理32\32\1字节数据，需要3次迭代可以完成，repeatTimes = 3；
168	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：32\32\1字节），这里的单位实际上是拼接后的方块矩阵的大小；	168	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：32\32\1字节），这里的单位实际上是拼接后的方块矩阵的大小；
169	- dstGap = 1，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为1（单位：512字节）；	169	- dstGap = 1，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为1（单位：512字节）；
170	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为0（单位：512字节）。	170	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为0（单位：512字节）。
	@@ -173,7 +173,7 @@ Kirin X90，支持数据类型：int8_t、half。
173		173
174	如下图示例：	174	如下图示例：
175		175
176	- - repeatTime和srcStride的解释和上图示例一致；	176	+ - repeatTimes和srcStride的解释和上图示例一致；
177	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；	177	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；
178	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为2（单位：512字节）。	178	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为2（单位：512字节）。
179		179
	@@ -181,7 +181,7 @@ Kirin X90，支持数据类型：int8_t、half。
181		181
182	- 对于half/bfloat16\_t数据类型，每次迭代处理16\16\2字节数据，可处理1个分形（一个分形512字节），每次迭代中，源操作数中1个16\*16分形将被转置。	182	- 对于half/bfloat16\_t数据类型，每次迭代处理16\16\2字节数据，可处理1个分形（一个分形512字节），每次迭代中，源操作数中1个16\*16分形将被转置。
183		183
184	- - 共需要处理1536字节的数据，每次迭代处理16\16\2字节数据，需要3次迭代可以完成，repeatTime = 3；	184	+ - 共需要处理1536字节的数据，每次迭代处理16\16\2字节数据，需要3次迭代可以完成，repeatTimes = 3；
185	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1 （单位：16\16\2字节）；	185	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1 （单位：16\16\2字节）；
186	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址无间隔；	186	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址无间隔；
187	- 该场景下，因为其分形即为方块矩阵，每个迭代处理一个分形，不存在迭代内分形的间隔，该参数设置无效。	187	- 该场景下，因为其分形即为方块矩阵，每个迭代处理一个分形，不存在迭代内分形的间隔，该参数设置无效。
	@@ -192,7 +192,7 @@ Kirin X90，支持数据类型：int8_t、half。
192		192
193	如下图示例：	193	如下图示例：
194		194
195	- - 共需要处理3072字节的数据，每次迭代处理16\16\4字节数据，需要3次迭代可以完成，repeatTime = 3；	195	+ - 共需要处理3072字节的数据，每次迭代处理16\16\4字节数据，需要3次迭代可以完成，repeatTimes = 3；
196	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：16\16\4字节），这里的单位实际上是拼接后的方块矩阵的大小；	196	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：16\16\4字节），这里的单位实际上是拼接后的方块矩阵的大小；
197	- dstGap = 1，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为1（单位：512字节）；	197	- dstGap = 1，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为1（单位：512字节）；
198	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形结束地址与后一个分形起始地址的间隔为0（单位：512字节）。	198	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形结束地址与后一个分形起始地址的间隔为0（单位：512字节）。
	@@ -201,7 +201,7 @@ Kirin X90，支持数据类型：int8_t、half。
201		201
202	如下图示例：	202	如下图示例：
203		203
204	- - repeatTime和srcStride的解释和上图示例一致；	204	+ - repeatTimes和srcStride的解释和上图示例一致；
205	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；	205	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；
206	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形结束地址与后一个分形起始地址的间隔为2（单位：512字节）。	206	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形结束地址与后一个分形起始地址的间隔为2（单位：512字节）。
207		207
	@@ -215,7 +215,7 @@ Kirin X90，支持数据类型：int8_t、half。
215		215
216	如下图示例：	216	如下图示例：
217		217
218	- - 共需要处理6144字节的数据，每次迭代处理64\64\0.5字节数据，需要3次迭代可以完成，repeatTime = 3；	218	+ - 共需要处理6144字节的数据，每次迭代处理64\64\0.5字节数据，需要3次迭代可以完成，repeatTimes = 3；
219	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：64\64\0.5字节），这里的单位实际上是拼接后的方块矩阵的大小；	219	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：64\64\0.5字节），这里的单位实际上是拼接后的方块矩阵的大小；
220	- dstGap = 1，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为1（单位：512字节）；	220	- dstGap = 1，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为1（单位：512字节）；
221	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为0（单位：512字节）。	221	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为0（单位：512字节）。
	@@ -224,7 +224,7 @@ Kirin X90，支持数据类型：int8_t、half。
224		224
225	如下图示例：	225	如下图示例：
226		226
227	- - repeatTime和srcStride的解释和上图示例一致；	227	+ - repeatTimes和srcStride的解释和上图示例一致；
228	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；	228	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；
229	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为2（单位：512字节）。	229	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为2（单位：512字节）。
230		230
	@@ -251,7 +251,7 @@ Kirin X90，支持数据类型：int8_t、half。
251		251
252	如下图示例：	252	如下图示例：
253		253
254	- - 共需要处理12288字节的数据，每次迭代处理64\64\0.5字节数据，需要3次迭代可以完成，repeatTime = 3；	254	+ - 共需要处理12288字节的数据，每次迭代处理64\64\0.5字节数据，需要3次迭代可以完成，repeatTimes = 3；
255	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：64\64\0.5字节）；	255	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：64\64\0.5字节）；
256	- dstGap = 7，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为7（单位：512字节）；	256	- dstGap = 7，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为7（单位：512字节）；
257	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为0（单位：512字节）；	257	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为0（单位：512字节）；
	@@ -261,7 +261,7 @@ Kirin X90，支持数据类型：int8_t、half。
261		261
262	如下图示例：	262	如下图示例：
263		263
264	- - repeatTime和srcStride的解释和上图示例一致；	264	+ - repeatTimes和srcStride的解释和上图示例一致；
265	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；	265	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；
266	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为2（单位：512字节）；	266	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形的结束地址与后一个分形起始地址的间隔为2（单位：512字节）；
267	- 外层for循环的循环次数是2，源操作数的地址偏移量是3（单位：64\64\0.5字节），目的操作数的地址偏移量是3（单位：64\64\0.5字节）。	267	- 外层for循环的循环次数是2，源操作数的地址偏移量是3（单位：64\64\0.5字节），目的操作数的地址偏移量是3（单位：64\64\0.5字节）。
	@@ -291,7 +291,7 @@ Kirin X90，支持数据类型：int8_t、half。
291		291
292	- 对于half数据类型，每次迭代处理16\16\2字节数据，可处理1个分形（一个分形512字节），每次迭代中，源操作数中1个16\*16分形将被转置。	292	- 对于half数据类型，每次迭代处理16\16\2字节数据，可处理1个分形（一个分形512字节），每次迭代中，源操作数中1个16\*16分形将被转置。
293		293
294	- - 共需要处理3072字节的数据，每次迭代处理16\16\2字节数据，需要3次迭代可以完成，repeatTime = 3；	294	+ - 共需要处理3072字节的数据，每次迭代处理16\16\2字节数据，需要3次迭代可以完成，repeatTimes = 3；
295	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1 （单位：16\16\2字节）；	295	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1 （单位：16\16\2字节）；
296	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址无间隔；	296	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址无间隔；
297	- 该场景下，因为其分形即为方块矩阵，每个迭代处理一个分形，不存在迭代内分形的间隔，dstFracGap参数设置无效。	297	- 该场景下，因为其分形即为方块矩阵，每个迭代处理一个分形，不存在迭代内分形的间隔，dstFracGap参数设置无效。
	@@ -303,7 +303,7 @@ Kirin X90，支持数据类型：int8_t、half。
303		303
304	如下图示例：	304	如下图示例：
305		305
306	- - 共需要处理6144字节的数据，每次迭代处理16\16\4字节数据，需要3次迭代可以完成，repeatTime = 3；	306	+ - 共需要处理6144字节的数据，每次迭代处理16\16\4字节数据，需要3次迭代可以完成，repeatTimes = 3；
307	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：16\16\4字节）；	307	- srcStride = 1，表示相邻迭代间，源操作数前一个方块矩阵与后一个方块矩阵起始地址的间隔为1（单位：16\16\4字节）；
308	- dstGap = 3，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为3（单位：512字节）；	308	- dstGap = 3，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址到下一个迭代第一个分形起始地址的间隔为3（单位：512字节）；
309	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形结束地址与后一个分形起始地址的间隔为0（单位：512字节）。	309	- dstFracGap = 0，表示每个迭代内目的操作数前一个分形结束地址与后一个分形起始地址的间隔为0（单位：512字节）。
	@@ -313,7 +313,7 @@ Kirin X90，支持数据类型：int8_t、half。
313		313
314	如下图示例：	314	如下图示例：
315		315
316	- - repeatTime和srcStride的解释和上图示例一致；	316	+ - repeatTimes和srcStride的解释和上图示例一致；
317	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；	317	- dstGap = 0，表示相邻迭代间，目的操作数前一个迭代第一个分形的结束地址和下一个迭代第一个分形起始地址无间隔；
318	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形结束地址与后一个分形起始地址的间隔为2（单位：512字节）；	318	- dstFracGap = 2，表示每个迭代内目的操作数前一个分形结束地址与后一个分形起始地址的间隔为2（单位：512字节）；
319	- 外层for循环的循环次数是2，源操作数的地址偏移量是3（单位：16\16\4字节），目的操作数的地址偏移量是3（单位：16\16\4字节）。	319	- 外层for循环的循环次数是2，源操作数的地址偏移量是3（单位：16\16\4字节），目的操作数的地址偏移量是3（单位：16\16\4字节）。

	@@ -98,7 +98,7 @@ DataCopy矩阵搬出接口支持多种随路能力的组合，需要设置不同
98	\| dstStride \| • 不开启NZ2ND功能（[NZ2NZ搬运](./关键特性说明/NZ2NZ.md)）：目的NZ矩阵中相邻Z排布的起始地址偏移，取值不为0，单位为datablock（32字节）。<br>• 开启随路[NZ2ND](./关键特性说明/NZ2ND.md)功能：目的ND矩阵每一行中的元素个数，取值不为0 ，单位为element。 \|	98	\| dstStride \| • 不开启NZ2ND功能（[NZ2NZ搬运](./关键特性说明/NZ2NZ.md)）：目的NZ矩阵中相邻Z排布的起始地址偏移，取值不为0，单位为datablock（32字节）。<br>• 开启随路[NZ2ND](./关键特性说明/NZ2ND.md)功能：目的ND矩阵每一行中的元素个数，取值不为0 ，单位为element。 \|
99	\| srcStride \| 源NZ矩阵中相邻Z排布的起始地址偏移，取值范围为srcStride∈[0, 65535]，单位为C0_Size（16*sizeof(T)），T为src的数据类型，其值应填成mSize对16向上取整。 \|	99	\| srcStride \| 源NZ矩阵中相邻Z排布的起始地址偏移，取值范围为srcStride∈[0, 65535]，单位为C0_Size（16*sizeof(T)），T为src的数据类型，其值应填成mSize对16向上取整。 \|
100	\| unitFlag \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag。<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位。<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>参数设置方案和特性细节可参考：Mmad计算中关键特性说明的[UnitFlag](../Mmad计算/关键特性说明/UnitFlag.md)章节。 \|	100	\| unitFlag \| unitFlag是一种Mmad指令和Fixpipe指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出。取值说明如下：<br>  • 0（2'b00）：不开启unitFlag。<br>  • 2（2'b10）：开启unitFlag，硬件执行完指令之后，不复位单元标记位。<br>  • 3（2'b11）：开启unitFlag，硬件执行完指令之后，复位单元标记位。<br>开启该功能时，须将Mmad指令和Fixpipe指令的unitFlag值设置为2或3。<br>参数设置方案和特性细节可参考：Mmad计算中关键特性说明的[UnitFlag](../Mmad计算/关键特性说明/UnitFlag.md)章节。 \|
101	-\| clipReluPre \| <!-- npu="950,A3,910b" id13 -->该参数仅在Atlas 200I/500 A2 推理产品支持。<br><!-- end id13 --><!-- npu="310b" id14 -->用于配置是否开启ClipReLU操作，参数类型为uint8_t，取值如下：0，不开启ClipReLU；1，开启ClipReLU，此时需要调用[SetFixPipeClipRelu](./寄存器配置说明/SetFixPipeClipRelu.md)来设置clipReLU的最大值。<br>  • 该操作在随路量化后进行，quantPre配置后才能使用，当前支持的量化模式有F322F16/DEQF16/VDEQF16/QF322B8_PRE/VQF322B8_PRE/REQ8/VREQ8。<br><!-- end id14 --> \|	101	+\| clipReluPre \| <!-- npu="950,A3,910b" id13 -->该参数仅在Atlas 200I/500 A2 推理产品支持。<br><!-- end id13 --><!-- npu="310b" id14 -->用于配置是否开启ClipReLU操作，参数类型为uint8_t，取值如下：0，不开启ClipReLU；1，开启ClipReLU，此时需要调用[SetFixPipeClipRelu](./寄存器配置说明/SetFixPipeClipRelu.md)来设置ClipReLU的最大值。<br>  • 该操作在随路量化后进行，quantPre配置后才能使用，当前支持的量化模式有F322F16/DEQF16/VDEQF16/QF322B8_PRE/VQF322B8_PRE/REQ8/VREQ8。<br><!-- end id14 --> \|
102	\| eltWiseOp \| <!-- npu="950,A3,910b" id15 -->该参数仅在Atlas 200I/500 A2 推理产品支持。<br><!-- end id15 --><!-- npu="310b" id16 -->用于配置是否开启Elementwise操作及操作模式。Elementwise操作是指进行随路量化后，可以逐个元素加/减一个LocalTensor，大小为mSize * nSize，具体LocalTensor地址相关参数需要调用[SetFixPipeAddr](./寄存器配置说明/SetFixPipeAddr.md)来设置。<br>eltWiseOp参数类型为uint8_t，取值如下：<br>  • 0：不开启Elementwise；<br>  • 1：Elementwise Addition；<br>  • 2：Elementwise Subtraction。<br><!-- end id16 --> \|	102	\| eltWiseOp \| <!-- npu="950,A3,910b" id15 -->该参数仅在Atlas 200I/500 A2 推理产品支持。<br><!-- end id15 --><!-- npu="310b" id16 -->用于配置是否开启Elementwise操作及操作模式。Elementwise操作是指进行随路量化后，可以逐个元素加/减一个LocalTensor，大小为mSize * nSize，具体LocalTensor地址相关参数需要调用[SetFixPipeAddr](./寄存器配置说明/SetFixPipeAddr.md)来设置。<br>eltWiseOp参数类型为uint8_t，取值如下：<br>  • 0：不开启Elementwise；<br>  • 1：Elementwise Addition；<br>  • 2：Elementwise Subtraction。<br><!-- end id16 --> \|
103	\| quantPre \| 用于控制量化模式，QuantMode_t类型，具体定义如下：<br>  • float/int32_t输出此需配置为QuantMode_t::NoQuant。<br>  • half/bfloat16_t输出，此参数需配置为QuantMode_t::F322F16/QuantMode_t::F322BF16。<br>  • 配置为scalar量化时，需要调用[SetFixpipePreQuantFlag](./寄存器配置说明/SetFixpipePreQuantFlag.md)接口来设置scalar量化参数。<br>  • 配置为tensor量化时，需要调用[SetFixPipeConfig](./寄存器配置说明/SetFixPipeConfig.md)来设置tensor量化参数，其中tensor量化参数需要通过DataCopy从L1 Buffer搬运至Fixpipe Buffer。<br>注：此参数需要用户手动配置，不会自动推导配置对应量化模式。<br><pre>enum QuantMode_t<br>{<br> NoQuant, // 不开启量化功能<br> F322F16, // Float32_2_Float16：float cast成half，cast mode为CAST_RINT模式<br> F322BF16, // Float32_2_BFloat16：float cast成bfloat16_t，cast mode为CAST_RINT模式<br> DEQF16, // DeQuant_Float16：int32_t量化成half，scalar量化<br> VDEQF16, // Vector_DeQuant_Float16：int32_t量化成half，tensor量化<br> QF322B8_PRE, // Quant_Float32_2_B8：float量化成int8_t/uint8_t，scalar量化<br> VQF322B8_PRE, // Vector_Quant_Float32_2_B8：float量化成int8_t/uint8_t，tensor量化<br> REQ8, // ReQuant_int8：int32_t量化成int8_t/uint8_t，scalar量化<br> VREQ8, // Vector_ReQuant_int8：int32_t量化成int8_t/uint8_t，tensor量化<br>};</pre> \|	103	\| quantPre \| 用于控制量化模式，QuantMode_t类型，具体定义如下：<br>  • float/int32_t输出此需配置为QuantMode_t::NoQuant。<br>  • half/bfloat16_t输出，此参数需配置为QuantMode_t::F322F16/QuantMode_t::F322BF16。<br>  • 配置为scalar量化时，需要调用[SetFixpipePreQuantFlag](./寄存器配置说明/SetFixpipePreQuantFlag.md)接口来设置scalar量化参数。<br>  • 配置为tensor量化时，需要调用[SetFixPipeConfig](./寄存器配置说明/SetFixPipeConfig.md)来设置tensor量化参数，其中tensor量化参数需要通过DataCopy从L1 Buffer搬运至Fixpipe Buffer。<br>注：此参数需要用户手动配置，不会自动推导配置对应量化模式。<br><pre>enum QuantMode_t<br>{<br> NoQuant, // 不开启量化功能<br> F322F16, // Float32_2_Float16：float cast成half，cast mode为CAST_RINT模式<br> F322BF16, // Float32_2_BFloat16：float cast成bfloat16_t，cast mode为CAST_RINT模式<br> DEQF16, // DeQuant_Float16：int32_t量化成half，scalar量化<br> VDEQF16, // Vector_DeQuant_Float16：int32_t量化成half，tensor量化<br> QF322B8_PRE, // Quant_Float32_2_B8：float量化成int8_t/uint8_t，scalar量化<br> VQF322B8_PRE, // Vector_Quant_Float32_2_B8：float量化成int8_t/uint8_t，tensor量化<br> REQ8, // ReQuant_int8：int32_t量化成int8_t/uint8_t，scalar量化<br> VREQ8, // Vector_ReQuant_int8：int32_t量化成int8_t/uint8_t，tensor量化<br>};</pre> \|
104	\| reluPre \| 用于配置ReLU操作的模式，类型为uint8_t，取值如下：<br>  • 0：不开启ReLU；<br>  • 1：Normal ReLU。 \|	104	\| reluPre \| 用于配置ReLU操作的模式，类型为uint8_t，取值如下：<br>  • 0：不开启ReLU；<br>  • 1：Normal ReLU。 \|
	@@ -375,7 +375,7 @@ DataCopy矩阵搬出接口支持多种随路能力的组合，需要设置不同
375	// AscendC::PipeBarrier<PIPE_FIX>();	375	// AscendC::PipeBarrier<PIPE_FIX>();
376	// // mov l0c to gm，量化操作后开启ClipReLU操作	376	// // mov l0c to gm，量化操作后开启ClipReLU操作
377	// intriParams.clipReluPre = 1;	377	// intriParams.clipReluPre = 1;
378	- // // 设置clip ReLU的值到寄存器	378	+ // // 设置ClipReLU的值到寄存器
379	// uint64_t clipReluVal = 0x3c00; // value 1, half	379	// uint64_t clipReluVal = 0x3c00; // value 1, half
380	// SetFixPipeClipRelu(clipReluVal);	380	// SetFixPipeClipRelu(clipReluVal);
381	// //mov l0c to gm，量化操作后，设置element-wise操作，Add	381	// //mov l0c to gm，量化操作后，设置element-wise操作，Add

	@@ -66,7 +66,7 @@ __aicore__ inline void SetFixPipeClipRelu(uint64_t config)
66		66
67	```cpp	67	```cpp
68	uint64_t clipReluMaxVal = 0x3c00;	68	uint64_t clipReluMaxVal = 0x3c00;
69	-SetFixPipeClipRelu(clipReluMaxVal); // 开启ReLU的情况下，先进行ReLU操作，再进行clip，clipReluMaxVal为通过该接口设置的最大值	69	+SetFixPipeClipRelu(clipReluMaxVal); // 开启ReLU的情况下，先进行ReLU操作，再进行ClipReLU，clipReluMaxVal为通过该接口设置的最大值
70	```	70	```
71		71
72	- 针对Atlas 200I/500 A2 推理产品，完整示例如下。随路格式转换数据搬运，通路：L0C Buffer-\>GM。	72	- 针对Atlas 200I/500 A2 推理产品，完整示例如下。随路格式转换数据搬运，通路：L0C Buffer-\>GM。
	@@ -222,7 +222,7 @@ SetFixPipeClipRelu(clipReluMaxVal); // 开启ReLU的情况下，先进行ReLU操
222	// AscendC::PipeBarrier<PIPE_FIX>();	222	// AscendC::PipeBarrier<PIPE_FIX>();
223	// // mov l0c to gm，量化操作后开启ClipReLU操作	223	// // mov l0c to gm，量化操作后开启ClipReLU操作
224	// intriParams.clipReluPre = 1;	224	// intriParams.clipReluPre = 1;
225	- // // 设置clip ReLU的值到寄存器	225	+ // // 设置ClipReLU的值到寄存器
226	// uint64_t clipReluVal = 0x3c00; // value 1, half	226	// uint64_t clipReluVal = 0x3c00; // value 1, half
227	// SetFixPipeClipRelu(clipReluVal);	227	// SetFixPipeClipRelu(clipReluVal);
228	// //mov l0c to gm，量化操作后，设置element-wise操作，Add	228	// //mov l0c to gm，量化操作后，设置element-wise操作，Add
	@@ -290,4 +290,4 @@ SetFixPipeClipRelu(clipReluMaxVal); // 开启ReLU的情况下，先进行ReLU操
290	}	290	}
291	KERNEL_CUBE_DATACOPY(half, int8_t, int8_t, int32_t, 128, 1, 1, QuantMode_t::DEQF16);	291	KERNEL_CUBE_DATACOPY(half, int8_t, int8_t, int32_t, 128, 1, 1, QuantMode_t::DEQF16);
292	```	292	```
293	-	293	+

	@@ -27,15 +27,15 @@ void UnPadTilingFunc(const ge::Shape srcShape, const uint32_t stackBufferSize, c
27	\| ascendcPlatform \| 输入 \| 传入硬件平台的信息，PlatformAscendC定义请参见[构造及析构函数](../../../Utils-API/平台信息获取/PlatformAscendC/构造及析构函数.md)。 \|	27	\| ascendcPlatform \| 输入 \| 传入硬件平台的信息，PlatformAscendC定义请参见[构造及析构函数](../../../Utils-API/平台信息获取/PlatformAscendC/构造及析构函数.md)。 \|
28	\| srcShape \| 输入 \| 输入Tensor的shape信息，shape为二维。 \|	28	\| srcShape \| 输入 \| 输入Tensor的shape信息，shape为二维。 \|
29	\| typeSize \| 输入 \| 输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。 \|	29	\| typeSize \| 输入 \| 输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。 \|
30	-\| maxValue \| 输出 \| UnPad接口能完成计算所需最大临时空间大小。<br><br>UnPad接口能完成计算所需的最大临时空间大小，超出该值的空间不会被该接口使用。在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。最大空间大小为0表示计算不需要临时空间。 <br>maxValue仅作为参考值，有可能大于Unified Buffer剩余空间的大小，该场景下，开发者需要根据Unified Buffer剩余空间的大小来选取合适的临时空间大小。 \|	30	+\| maxValue \| 输出 \| UnPad接口能完成计算所需的最大临时空间大小，超出该值的空间不会被该接口使用。在最小临时空间-最大临时空间范围内，随着临时空间增大，kernel侧接口计算性能会有一定程度的优化提升。为了达到更好的性能，开发者可以根据实际的内存使用情况进行空间预留/申请。最大空间大小为0表示计算不需要临时空间。<br>maxValue仅作为参考值，有可能大于Unified Buffer剩余空间的大小，该场景下，开发者需要根据Unified Buffer剩余空间的大小来选取合适的临时空间大小。 \|
31	-\| minValue \| 输出 \| UnPad接口能完成计算所需最小临时空间大小。<br><br>Pad接口能完成计算所需最小临时空间大小。为保证功能正确，接口计算时预留/申请的临时空间不能小于该数值。最小空间大小为0表示计算不需要临时空间。 \|	31	+\| minValue \| 输出 \| UnPad接口能完成计算所需最小临时空间大小。为保证功能正确，接口计算时预留/申请的临时空间不能小于该数值。最小空间大小为0表示计算不需要临时空间。 \|
32		32
33	表2 UnPadTilingFunc接口参数说明	33	表2 UnPadTilingFunc接口参数说明
34		34
35	\| 参数名 \| 输入/输出 \| 含义 \|	35	\| 参数名 \| 输入/输出 \| 含义 \|
36	\| --- \| --- \| --- \|	36	\| --- \| --- \| --- \|
37	\| srcShape \| 输入 \| 输入Tensor的shape信息，shape为二维。 \|	37	\| srcShape \| 输入 \| 输入Tensor的shape信息，shape为二维。 \|
38	-\| stackBufferSize \| 输入 \| 可供UnPad接口计算的空间大小。 \|	38	+\| stackBufferSize \| 输入 \| 可供UnPad接口计算的临时空间大小，单位为Byte。 \|
39	\| typeSize \| 输入 \| 输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。 \|	39	\| typeSize \| 输入 \| 输入的数据类型大小，单位为字节。比如输入的数据类型为half，此处应传入2。 \|
40	\| tiling \| 输出 \| 输出UnPad接口所需的tiling信息。 \|	40	\| tiling \| 输出 \| 输出UnPad接口所需的tiling信息。 \|
41		41
	@@ -49,7 +49,7 @@ void UnPadTilingFunc(const ge::Shape srcShape, const uint32_t stackBufferSize, c
49		49
50	## 调用示例	50	## 调用示例
51		51
52	-如下样例介绍了使用UnPad高阶API时host侧获取Tiling参数的流程以及该参数如何在kernel侧使用。样例中原始shape的大小为\[320, 64\]，需要unpad的目标shape大小为\[320, 63\]，输入的数据类型为half。	52	+如下样例介绍了使用UnPad高阶API时host侧获取Tiling参数的流程以及该参数如何在kernel侧使用。样例中原始shape的大小为\[320, 64\]，unpad后的目标shape大小为\[320, 63\]，输入的数据类型为half。
53		53
54	1. 将UnPadTiling结构体参数增加至TilingData结构体，作为TilingData结构体的一个字段。	54	1. 将UnPadTiling结构体参数增加至TilingData结构体，作为TilingData结构体的一个字段。
55		55
	@@ -62,7 +62,7 @@ void UnPadTilingFunc(const ge::Shape srcShape, const uint32_t stackBufferSize, c
62	END_TILING_DATA_DEF;	62	END_TILING_DATA_DEF;
63	```	63	```
64		64
65	-2. Tiling实现函数中，首先调用GetUnPadMaxMinTmpSize接口获取UnPad接口能完成计算所需最大/最小临时空间大小，根据该范围结合实际的内存使用情况设置合适的空间大小；然后根据输入shape、剩余的可供计算的空间大小等信息获取UnPad kernel侧接口所需tiling参数。	65	+2. Tiling实现函数中，首先调用GetUnPadMaxMinTmpSize接口获取UnPad接口能完成计算所需最大和最小临时空间大小，根据该范围结合实际的内存使用情况设置合适的空间大小；然后根据输入shape、剩余的可供计算的空间大小等信息获取UnPad kernel侧接口所需tiling参数。
66		66
67	```	67	```
68	namespace optiling {	68	namespace optiling {
	@@ -84,7 +84,7 @@ void UnPadTilingFunc(const ge::Shape srcShape, const uint32_t stackBufferSize, c
84	auto platformInfo = context->GetPlatformInfo();	84	auto platformInfo = context->GetPlatformInfo();
85	auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);	85	auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
86	AscendC::GetUnPadMaxMinTmpSize(ascendcPlatform, srcShape, sizeof(half), maxValue, minValue);	86	AscendC::GetUnPadMaxMinTmpSize(ascendcPlatform, srcShape, sizeof(half), maxValue, minValue);
87	- // 本样例中仅作为样例说明，获取最小值并传入，来保证功能正确，开发者可以根据需要传入合适的空间大小	87	+ // 本样例仅用于说明，获取最小值并传入，来保证功能正确，开发者可以根据需要传入合适的空间大小
88	const uint32_t localWorkSpaceSize = minValue;	88	const uint32_t localWorkSpaceSize = minValue;
89	AscendC::UnPadTilingFunc(srcShape, localWorkSpaceSize , sizeof(half), tiling.unpadTilingData);	89	AscendC::UnPadTilingFunc(srcShape, localWorkSpaceSize , sizeof(half), tiling.unpadTilingData);
90	...	90	...

	@@ -42,26 +42,26 @@ template <typename T> __aicore__ inline GemmTiling GetGemmTiling(uint32_t m, uin
42	* @param [in] m Number of rows of src0	42	* @param [in] m Number of rows of src0
43	* @param [in] n Number of rows of src1	43	* @param [in] n Number of rows of src1
44	* @param [in] k Number of columns of src1	44	* @param [in] k Number of columns of src1
45	- * @param [in] tilling.blockSize size of block	45	+ * @param [in] tiling.blockSize size of block
46	- * @param [in] tilling.mNum args of m	46	+ * @param [in] tiling.mNum args of m
47	- * @param [in] tilling.nNum args of n	47	+ * @param [in] tiling.nNum args of n
48	- * @param [in] tilling.kNum args of k	48	+ * @param [in] tiling.kNum args of k
49	- * @param [in] tilling.roundM/N/K Rounding parameter	49	+ * @param [in] tiling.roundM/N/K Rounding parameter
50	- * @param [in] tilling.c0Size The byte length of a block	50	+ * @param [in] tiling.c0Size The byte length of a block
51	- * @param [in] tilling.dtypeSize Byte length of the incoming data type	51	+ * @param [in] tiling.dtypeSize Byte length of the incoming data type
52	- * @param [in] tilling.m/n/kBlockNum Number of blocks of m/n/k axis	52	+ * @param [in] tiling.m/n/kBlockNum Number of blocks of m/n/k axis
53	- * @param [in] tilling.m/n/kIterNum Number of traversal dimensions	53	+ * @param [in] tiling.m/n/kIterNum Number of traversal dimensions
54	- * @param [in] tilling.m/k/nTileBlock Number of M/N/K axis cutting blocks	54	+ * @param [in] tiling.m/k/nTileBlock Number of M/N/K axis cutting blocks
55	- * @param [in] tilling.m/n/kHasTailNumber of tail blocks of M/K/N axis	55	+ * @param [in] tiling.m/n/kHasTailNumber of tail blocks of M/K/N axis
56	- * @param [in] tilling.kHasTileEle Judge whether the tail block exists	56	+ * @param [in] tiling.kHasTileEle Judge whether the tail block exists
57	- * @param [in] tilling.KtailEle K-axis tail block element	57	+ * @param [in] tiling.KtailEle K-axis tail block element
58	- * @param [in] tilling.kThreadNum K-axis passes	58	+ * @param [in] tiling.kThreadNum K-axis passes
59	* @param [in] partialsum judge whether the calculation result is moved out	59	* @param [in] partialsum judge whether the calculation result is moved out
60	* @param [in] initValue Initialization parameters	60	* @param [in] initValue Initialization parameters
61	*/	61	*/
62	template <typename T, typename U, typename S>	62	template <typename T, typename U, typename S>
63	__aicore__ inline __inout_pipe__(V) void Gemm(const LocalTensor<T>& dst, const LocalTensor<U>& src0,	63	__aicore__ inline __inout_pipe__(V) void Gemm(const LocalTensor<T>& dst, const LocalTensor<U>& src0,
64	- const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling tilling,	64	+ const LocalTensor<S>& src1, const uint32_t m, const uint32_t k, const uint32_t n, GemmTiling tiling,
65	bool partialsum = true, int32_t initValue = 0);	65	bool partialsum = true, int32_t initValue = 0);
66	} // namespace AscendC	66	} // namespace AscendC
67		67