已合并
docs: 锚点与相邻文字合并至同一行 #6592
ai_xin创建于 28 天前
docs: 锚点与相邻文字合并至同一行 #6592
已合并
ai_xin创建于 28 天前
共 57 个文件变更+117-262
@@ -6,9 +6,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
6 6 
7- 支持的数据类型有变化。7- 支持的数据类型有变化。
8 8 
9- **表1** 数据类型兼容性情况9+ **表1** 数据类型兼容性情况<a name="table1534819819338"></a>
10 10 
11- <a name="table1534819819338"></a>
12 <table><thead align="left"><tr id="row1234720819332"><th class="cellrowborder" valign="top" width="20.849999999999998%" id="mcps1.2.6.1.1"><p id="p43471486333"><a name="p43471486333"></a><a name="p43471486333"></a>A矩阵</p>11 <table><thead align="left"><tr id="row1234720819332"><th class="cellrowborder" valign="top" width="20.849999999999998%" id="mcps1.2.6.1.1"><p id="p43471486333"><a name="p43471486333"></a><a name="p43471486333"></a>A矩阵</p>
13 </th>12 </th>
14 <th class="cellrowborder" valign="top" width="19.05%" id="mcps1.2.6.1.2"><p id="p14347686337"><a name="p14347686337"></a><a name="p14347686337"></a>B矩阵</p>13 <th class="cellrowborder" valign="top" width="19.05%" id="mcps1.2.6.1.2"><p id="p14347686337"><a name="p14347686337"></a><a name="p14347686337"></a>B矩阵</p>
@@ -105,9 +104,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
105 104 
106## 其它高阶API<a name="section14195124865717"></a>105## 其它高阶API<a name="section14195124865717"></a>
107 106 
108-**表2** 数学计算107+**表2** 数学计算<a name="table1011733113422"></a>
109 108 
110-<a name="table1011733113422"></a>
111<table><thead align="left"><tr id="row3118531154217"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p0118031104211"><a name="p0118031104211"></a><a name="p0118031104211"></a>AscendC高阶API</p>109<table><thead align="left"><tr id="row3118531154217"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p0118031104211"><a name="p0118031104211"></a><a name="p0118031104211"></a>AscendC高阶API</p>
112</th>110</th>
113<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p8118143124210"><a name="p8118143124210"></a><a name="p8118143124210"></a>兼容说明</p>111<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p8118143124210"><a name="p8118143124210"></a><a name="p8118143124210"></a>兼容说明</p>
@@ -142,9 +140,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
142</tbody>140</tbody>
143</table>141</table>
144 142 
145-**表3** 激活函数143+**表3** 激活函数<a name="table192231726173811"></a>
146 144 
147-<a name="table192231726173811"></a>
148<table><thead align="left"><tr id="row1422362623814"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p1041413536446"><a name="p1041413536446"></a><a name="p1041413536446"></a>AscendC高阶API</p>145<table><thead align="left"><tr id="row1422362623814"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p1041413536446"><a name="p1041413536446"></a><a name="p1041413536446"></a>AscendC高阶API</p>
149</th>146</th>
150<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p9414175314412"><a name="p9414175314412"></a><a name="p9414175314412"></a>兼容说明</p>147<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p9414175314412"><a name="p9414175314412"></a><a name="p9414175314412"></a>兼容说明</p>
@@ -159,9 +156,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
159</tbody>156</tbody>
160</table>157</table>
161 158 
162-**表4** 数据归一化159+**表4** 数据归一化<a name="table1488319115396"></a>
163 160 
164-<a name="table1488319115396"></a>
165<table><thead align="left"><tr id="row138831103911"><th class="cellrowborder" valign="top" width="49.980000000000004%" id="mcps1.2.3.1.1"><p id="p8665135824416"><a name="p8665135824416"></a><a name="p8665135824416"></a>AscendC高阶API</p>161<table><thead align="left"><tr id="row138831103911"><th class="cellrowborder" valign="top" width="49.980000000000004%" id="mcps1.2.3.1.1"><p id="p8665135824416"><a name="p8665135824416"></a><a name="p8665135824416"></a>AscendC高阶API</p>
166</th>162</th>
167<th class="cellrowborder" valign="top" width="50.019999999999996%" id="mcps1.2.3.1.2"><p id="p1466513582447"><a name="p1466513582447"></a><a name="p1466513582447"></a>兼容说明</p>163<th class="cellrowborder" valign="top" width="50.019999999999996%" id="mcps1.2.3.1.2"><p id="p1466513582447"><a name="p1466513582447"></a><a name="p1466513582447"></a>兼容说明</p>
@@ -194,9 +190,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
194</tbody>190</tbody>
195</table>191</table>
196 192 
197-**表5** 量化操作193+**表5** 量化操作<a name="table432310426394"></a>
198 194 
199-<a name="table432310426394"></a>
200<table><thead align="left"><tr id="row18324164212392"><th class="cellrowborder" valign="top" width="33.95%" id="mcps1.2.3.1.1"><p id="p1247120164516"><a name="p1247120164516"></a><a name="p1247120164516"></a>AscendC高阶API</p>195<table><thead align="left"><tr id="row18324164212392"><th class="cellrowborder" valign="top" width="33.95%" id="mcps1.2.3.1.1"><p id="p1247120164516"><a name="p1247120164516"></a><a name="p1247120164516"></a>AscendC高阶API</p>
201</th>196</th>
202<th class="cellrowborder" valign="top" width="66.05%" id="mcps1.2.3.1.2"><p id="p1247130184515"><a name="p1247130184515"></a><a name="p1247130184515"></a>兼容说明</p>197<th class="cellrowborder" valign="top" width="66.05%" id="mcps1.2.3.1.2"><p id="p1247130184515"><a name="p1247130184515"></a><a name="p1247130184515"></a>兼容说明</p>
@@ -230,9 +225,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
230</tbody>225</tbody>
231</table>226</table>
232 227 
233-**表6** 归约操作228+**表6** 归约操作<a name="table17561457193915"></a>
234 229 
235-<a name="table17561457193915"></a>
236<table><thead align="left"><tr id="row8756357163920"><th class="cellrowborder" valign="top" width="60.709999999999994%" id="mcps1.2.3.1.1"><p id="p102601213455"><a name="p102601213455"></a><a name="p102601213455"></a>AscendC高阶API</p>230<table><thead align="left"><tr id="row8756357163920"><th class="cellrowborder" valign="top" width="60.709999999999994%" id="mcps1.2.3.1.1"><p id="p102601213455"><a name="p102601213455"></a><a name="p102601213455"></a>AscendC高阶API</p>
237</th>231</th>
238<th class="cellrowborder" valign="top" width="39.290000000000006%" id="mcps1.2.3.1.2"><p id="p726010113457"><a name="p726010113457"></a><a name="p726010113457"></a>兼容说明</p>232<th class="cellrowborder" valign="top" width="39.290000000000006%" id="mcps1.2.3.1.2"><p id="p726010113457"><a name="p726010113457"></a><a name="p726010113457"></a>兼容说明</p>
@@ -259,9 +253,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
259</tbody>253</tbody>
260</table>254</table>
261 255 
262-**表7** 排序操作256+**表7** 排序操作<a name="table4855685408"></a>
263 257 
264-<a name="table4855685408"></a>
265<table><thead align="left"><tr id="row78561987407"><th class="cellrowborder" valign="top" width="50.33%" id="mcps1.2.3.1.1"><p id="p18599528455"><a name="p18599528455"></a><a name="p18599528455"></a>AscendC高阶API</p>258<table><thead align="left"><tr id="row78561987407"><th class="cellrowborder" valign="top" width="50.33%" id="mcps1.2.3.1.1"><p id="p18599528455"><a name="p18599528455"></a><a name="p18599528455"></a>AscendC高阶API</p>
266</th>259</th>
267<th class="cellrowborder" valign="top" width="49.669999999999995%" id="mcps1.2.3.1.2"><p id="p115995284510"><a name="p115995284510"></a><a name="p115995284510"></a>兼容说明</p>260<th class="cellrowborder" valign="top" width="49.669999999999995%" id="mcps1.2.3.1.2"><p id="p115995284510"><a name="p115995284510"></a><a name="p115995284510"></a>兼容说明</p>
@@ -288,9 +281,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
288</tbody>281</tbody>
289</table>282</table>
290 283 
291-**表8** 索引计算284+**表8** 索引计算<a name="table748063194016"></a>
292 285 
293-<a name="table748063194016"></a>
294<table><thead align="left"><tr id="row8480531134011"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p1944815184518"><a name="p1944815184518"></a><a name="p1944815184518"></a>AscendC高阶API</p>286<table><thead align="left"><tr id="row8480531134011"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p1944815184518"><a name="p1944815184518"></a><a name="p1944815184518"></a>AscendC高阶API</p>
295</th>287</th>
296<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p1244845124510"><a name="p1244845124510"></a><a name="p1244845124510"></a>兼容说明</p>288<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p1244845124510"><a name="p1244845124510"></a><a name="p1244845124510"></a>兼容说明</p>
@@ -306,9 +298,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
306</tbody>298</tbody>
307</table>299</table>
308 300 
309-**表9** 数据过滤301+**表9** 数据过滤<a name="table1937505494013"></a>
310 302 
311-<a name="table1937505494013"></a>
312<table><thead align="left"><tr id="row03769549406"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p182236817450"><a name="p182236817450"></a><a name="p182236817450"></a>AscendC高阶API</p>303<table><thead align="left"><tr id="row03769549406"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p182236817450"><a name="p182236817450"></a><a name="p182236817450"></a>AscendC高阶API</p>
313</th>304</th>
314<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p422388174510"><a name="p422388174510"></a><a name="p422388174510"></a>兼容说明</p>305<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p422388174510"><a name="p422388174510"></a><a name="p422388174510"></a>兼容说明</p>
@@ -329,9 +320,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
329</tbody>320</tbody>
330</table>321</table>
331 322 
332-**表10** 张量变换323+**表10** 张量变换<a name="table68071919174115"></a>
333 324 
334-<a name="table68071919174115"></a>
335<table><thead align="left"><tr id="row1880713195411"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p2655161904519"><a name="p2655161904519"></a><a name="p2655161904519"></a>AscendC高阶API</p>325<table><thead align="left"><tr id="row1880713195411"><th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.1"><p id="p2655161904519"><a name="p2655161904519"></a><a name="p2655161904519"></a>AscendC高阶API</p>
336</th>326</th>
337<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p14655719124519"><a name="p14655719124519"></a><a name="p14655719124519"></a>兼容说明</p>327<th class="cellrowborder" valign="top" width="50%" id="mcps1.2.3.1.2"><p id="p14655719124519"><a name="p14655719124519"></a><a name="p14655719124519"></a>兼容说明</p>
@@ -366,9 +356,8 @@ Ascend C高阶API基本兼容[NPU架构版本3510](../../../programming_guide/la
366</tbody>356</tbody>
367</table>357</table>
368 358 
369-**表11** Hccl359+**表11** Hccl<a name="table12265143614414"></a>
370 360 
371-<a name="table12265143614414"></a>
372<table><thead align="left"><tr id="row1126583611410"><th class="cellrowborder" valign="top" width="43.35%" id="mcps1.2.3.1.1"><p id="p16138821104518"><a name="p16138821104518"></a><a name="p16138821104518"></a>AscendC高阶API</p>361<table><thead align="left"><tr id="row1126583611410"><th class="cellrowborder" valign="top" width="43.35%" id="mcps1.2.3.1.1"><p id="p16138821104518"><a name="p16138821104518"></a><a name="p16138821104518"></a>AscendC高阶API</p>
373</th>362</th>
374<th class="cellrowborder" valign="top" width="56.65%" id="mcps1.2.3.1.2"><p id="p61381215456"><a name="p61381215456"></a><a name="p61381215456"></a>兼容说明</p>363<th class="cellrowborder" valign="top" width="56.65%" id="mcps1.2.3.1.2"><p id="p61381215456"><a name="p61381215456"></a><a name="p61381215456"></a>兼容说明</p>
@@ -8,9 +8,8 @@
8![](../figures/api_hierarchy.png "Ascend-C-API层次结构")8![](../figures/api_hierarchy.png "Ascend-C-API层次结构")
9 9 
10 10 
11-**表1** Ascend C API兼容策略11+**表1** Ascend C API兼容策略<a name="table18109729593"></a>
12 12 
13-<a name="table18109729593"></a>
14<table><thead align="left"><tr id="row810952912919"><th class="cellrowborder" valign="top" width="18.01%" id="mcps1.2.3.1.1"><p id="p141091429991"><a name="p141091429991"></a><a name="p141091429991"></a>API层级</p>13<table><thead align="left"><tr id="row810952912919"><th class="cellrowborder" valign="top" width="18.01%" id="mcps1.2.3.1.1"><p id="p141091429991"><a name="p141091429991"></a><a name="p141091429991"></a>API层级</p>
15</th>14</th>
16<th class="cellrowborder" valign="top" width="81.99%" id="mcps1.2.3.1.2"><p id="p0109162915915"><a name="p0109162915915"></a><a name="p0109162915915"></a>兼容策略</p>15<th class="cellrowborder" valign="top" width="81.99%" id="mcps1.2.3.1.2"><p id="p0109162915915"><a name="p0109162915915"></a><a name="p0109162915915"></a>兼容策略</p>
@@ -16,9 +16,8 @@ for i in range(g):
16 16 
17优化分析以如下算子规格为例:17优化分析以如下算子规格为例:
18 18 
19-**表1** 算子规格19+**表1** 算子规格<a name="table162628535334"></a>
20 20 
21-<a name="table162628535334"></a>
22<table><thead align="left"><tr id="row17263115333315"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p528831925514"><a name="p528831925514"></a><a name="p528831925514"></a>input</p>21<table><thead align="left"><tr id="row17263115333315"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p528831925514"><a name="p528831925514"></a><a name="p528831925514"></a>input</p>
23</th>22</th>
24<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1726365314337"><a name="p1726365314337"></a><a name="p1726365314337"></a>shape</p>23<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1726365314337"><a name="p1726365314337"></a><a name="p1726365314337"></a>shape</p>
@@ -10,9 +10,8 @@
10 10 
11本案例的算子规格如下:11本案例的算子规格如下:
12 12 
13-**表1** 算子规格13+**表1** 算子规格<a name="table101751125175213"></a>
14 14 
15-<a name="table101751125175213"></a>
16<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>15<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>
17</th>16</th>
18<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>17<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>
@@ -113,9 +112,8 @@
113- 优化后的Cube流水图如下所示,不使用随路ND2NZ指令对B矩阵进行格式转换后,MTE2的占比明显下降。![](../../../figures/zh-cn_image_0000002443251985.png)112- 优化后的Cube流水图如下所示,不使用随路ND2NZ指令对B矩阵进行格式转换后,MTE2的占比明显下降。![](../../../figures/zh-cn_image_0000002443251985.png)
114- 优化后的Profiling数据如下,可以看到同时使用Cube单元和Vector单元,aic\_time最大耗时90.95us,其中aic\_mte2\_ratio占比明显降低。![](../../../figures/zh-cn_image_0000002381162386.png)![](../../../figures/zh-cn_image_0000002414561861.png)113- 优化后的Profiling数据如下,可以看到同时使用Cube单元和Vector单元,aic\_time最大耗时90.95us,其中aic\_mte2\_ratio占比明显降低。![](../../../figures/zh-cn_image_0000002381162386.png)![](../../../figures/zh-cn_image_0000002414561861.png)
115 114 
116-**表2** 端到端性能对比115+**表2** 端到端性能对比<a name="table135810433437"></a>
117 116 
118-<a name="table135810433437"></a>
119<table><thead align="left"><tr id="row65820433437"><th class="cellrowborder" valign="top" width="22.61%" id="mcps1.2.5.1.1"><p id="p75827438435"><a name="p75827438435"></a><a name="p75827438435"></a>优化方法</p>117<table><thead align="left"><tr id="row65820433437"><th class="cellrowborder" valign="top" width="22.61%" id="mcps1.2.5.1.1"><p id="p75827438435"><a name="p75827438435"></a><a name="p75827438435"></a>优化方法</p>
120</th>118</th>
121<th class="cellrowborder" valign="top" width="21.88%" id="mcps1.2.5.1.2"><p id="p3582204311439"><a name="p3582204311439"></a><a name="p3582204311439"></a>总耗时(us)</p>119<th class="cellrowborder" valign="top" width="21.88%" id="mcps1.2.5.1.2"><p id="p3582204311439"><a name="p3582204311439"></a><a name="p3582204311439"></a>总耗时(us)</p>
@@ -11,9 +11,8 @@
11 11 
12本案例的算子规格如下:12本案例的算子规格如下:
13 13 
14-**表1** 算子规格14+**表1** 算子规格<a name="table568792363119"></a>
15 15 
16-<a name="table568792363119"></a>
17<table><thead align="left"><tr id="row1368792319318"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p186887235312"><a name="p186887235312"></a><a name="p186887235312"></a>输入</p>16<table><thead align="left"><tr id="row1368792319318"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p186887235312"><a name="p186887235312"></a><a name="p186887235312"></a>输入</p>
18</th>17</th>
19<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1268862303114"><a name="p1268862303114"></a><a name="p1268862303114"></a>Shape</p>18<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1268862303114"><a name="p1268862303114"></a><a name="p1268862303114"></a>Shape</p>
@@ -14,9 +14,8 @@
14 14 
15本案例的算子规格如下:15本案例的算子规格如下:
16 16 
17-**表1** 算子规格17+**表1** 算子规格<a name="table101751125175213"></a>
18 18 
19-<a name="table101751125175213"></a>
20<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>19<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>
21</th>20</th>
22<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>21<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>
@@ -12,9 +12,8 @@
12 12 
13本案例的算子规格如下:13本案例的算子规格如下:
14 14 
15-**表1** 算子规格15+**表1** 算子规格<a name="table101751125175213"></a>
16 16 
17-<a name="table101751125175213"></a>
18<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>17<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>
19</th>18</th>
20<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>19<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>
@@ -14,9 +14,8 @@
14 14 
15本案例的算子规格如下:15本案例的算子规格如下:
16 16 
17-**表1** 算子规格17+**表1** 算子规格<a name="table9129161439"></a>
18 18 
19-<a name="table9129161439"></a>
20<table><thead align="left"><tr id="row4121116114317"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1712111610439"><a name="p1712111610439"></a><a name="p1712111610439"></a>输入</p>19<table><thead align="left"><tr id="row4121116114317"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1712111610439"><a name="p1712111610439"></a><a name="p1712111610439"></a>输入</p>
21</th>20</th>
22<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p15129162437"><a name="p15129162437"></a><a name="p15129162437"></a>Shape</p>21<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p15129162437"><a name="p15129162437"></a><a name="p15129162437"></a>Shape</p>
@@ -17,9 +17,8 @@
17 17 
18本案例的算子规格如下:18本案例的算子规格如下:
19 19 
20-**表1** 算子规格20+**表1** 算子规格<a name="table101751125175213"></a>
21 21 
22-<a name="table101751125175213"></a>
23<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>22<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>
24</th>23</th>
25<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>24<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>
@@ -22,9 +22,8 @@
22 22 
23本案例的算子规格如下:23本案例的算子规格如下:
24 24 
25-**表1** 算子用例规格25+**表1** 算子用例规格<a name="table15465191317123"></a>
26 26 
27-<a name="table15465191317123"></a>
28<table><thead align="left"><tr id="row184651013131217"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p24653132122"><a name="p24653132122"></a><a name="p24653132122"></a>输入</p>27<table><thead align="left"><tr id="row184651013131217"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p24653132122"><a name="p24653132122"></a><a name="p24653132122"></a>输入</p>
29</th>28</th>
30<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p13465111311213"><a name="p13465111311213"></a><a name="p13465111311213"></a>Shape</p>29<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p13465111311213"><a name="p13465111311213"></a><a name="p13465111311213"></a>Shape</p>
@@ -15,9 +15,8 @@
15 15 
16本案例的算子规格如下:16本案例的算子规格如下:
17 17 
18-**表1** 算子规格18+**表1** 算子规格<a name="table101751125175213"></a>
19 19 
20-<a name="table101751125175213"></a>
21<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>20<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>
22</th>21</th>
23<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>22<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>
@@ -18,9 +18,8 @@
18 18 
19本案例的算子规格如下:19本案例的算子规格如下:
20 20 
21-**表1** 算子规格21+**表1** 算子规格<a name="table101751125175213"></a>
22 22 
23-<a name="table101751125175213"></a>
24<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>23<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>
25</th>24</th>
26<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>25<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>
@@ -4,9 +4,8 @@
4 4 
5- Tiling优化5- Tiling优化
6 6 
7- **表1** Tiling优化策略总览7+ **表1** Tiling优化策略总览<a name="table11377439144919"></a>
8 8 
9- <a name="table11377439144919"></a>
10 <table><thead align="left"><tr id="row7377639104916"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p4377123917495"><a name="p4377123917495"></a><a name="p4377123917495"></a>分类</p>9 <table><thead align="left"><tr id="row7377639104916"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p4377123917495"><a name="p4377123917495"></a><a name="p4377123917495"></a>分类</p>
11 </th>10 </th>
12 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p2377193920493"><a name="p2377193920493"></a><a name="p2377193920493"></a>适用场景</p>11 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p2377193920493"><a name="p2377193920493"></a><a name="p2377193920493"></a>适用场景</p>
@@ -27,9 +26,8 @@
27 26 
28- 并行度优化27- 并行度优化
29 28 
30- **表2** 并行度优化策略总览29+ **表2** 并行度优化策略总览<a name="table225712267501"></a>
31 30 
32- <a name="table225712267501"></a>
33 <table><thead align="left"><tr id="row18257182614507"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p1925717267508"><a name="p1925717267508"></a><a name="p1925717267508"></a>分类</p>31 <table><thead align="left"><tr id="row18257182614507"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p1925717267508"><a name="p1925717267508"></a><a name="p1925717267508"></a>分类</p>
34 </th>32 </th>
35 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p182571626155016"><a name="p182571626155016"></a><a name="p182571626155016"></a>适用场景</p>33 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p182571626155016"><a name="p182571626155016"></a><a name="p182571626155016"></a>适用场景</p>
@@ -65,9 +63,8 @@
65 63 
66- 内存优化64- 内存优化
67 65 
68- **表3** 内存优化策略总览66+ **表3** 内存优化策略总览<a name="table136011854115018"></a>
69 67 
70- <a name="table136011854115018"></a>
71 <table><thead align="left"><tr id="row1060112547508"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p860113549502"><a name="p860113549502"></a><a name="p860113549502"></a>分类</p>68 <table><thead align="left"><tr id="row1060112547508"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p860113549502"><a name="p860113549502"></a><a name="p860113549502"></a>分类</p>
72 </th>69 </th>
73 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p460185475011"><a name="p460185475011"></a><a name="p460185475011"></a>适用场景</p>70 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p460185475011"><a name="p460185475011"></a><a name="p460185475011"></a>适用场景</p>
@@ -96,9 +93,8 @@
96 93 
97- Scalar优化94- Scalar优化
98 95 
99- **表4** Scalar优化策略总览96+ **表4** Scalar优化策略总览<a name="table214614239517"></a>
100 97 
101- <a name="table214614239517"></a>
102 <table><thead align="left"><tr id="row514716230519"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p16147112355116"><a name="p16147112355116"></a><a name="p16147112355116"></a>分类</p>98 <table><thead align="left"><tr id="row514716230519"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p16147112355116"><a name="p16147112355116"></a><a name="p16147112355116"></a>分类</p>
103 </th>99 </th>
104 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p12147112318519"><a name="p12147112318519"></a><a name="p12147112318519"></a>适用场景</p>100 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p12147112318519"><a name="p12147112318519"></a><a name="p12147112318519"></a>适用场景</p>
@@ -127,9 +123,8 @@
127 123 
128- 搬运优化124- 搬运优化
129 125 
130- **表5** 搬运优化策略总览126+ **表5** 搬运优化策略总览<a name="table2401471518"></a>
131 127 
132- <a name="table2401471518"></a>
133 <table><thead align="left"><tr id="row44094713510"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p114013470517"><a name="p114013470517"></a><a name="p114013470517"></a>分类</p>128 <table><thead align="left"><tr id="row44094713510"><th class="cellrowborder" valign="top" width="30.086991300869915%" id="mcps1.2.4.1.1"><p id="p114013470517"><a name="p114013470517"></a><a name="p114013470517"></a>分类</p>
134 </th>129 </th>
135 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p164015478513"><a name="p164015478513"></a><a name="p164015478513"></a>适用场景</p>130 <th class="cellrowborder" valign="top" width="40.82591740825918%" id="mcps1.2.4.1.2"><p id="p164015478513"><a name="p164015478513"></a><a name="p164015478513"></a>适用场景</p>
@@ -11,9 +11,8 @@
11 11 
12 非融合算子,只有矩阵计算的场景。即相较于MIX模式(包含矩阵计算和矢量计算),没有矢量计算的场景。本案例的算子规格如下:12 非融合算子,只有矩阵计算的场景。即相较于MIX模式(包含矩阵计算和矢量计算),没有矢量计算的场景。本案例的算子规格如下:
13 13 
14-**表1** 算子用例规格14+**表1** 算子用例规格<a name="table15465191317123"></a>
15 15 
16-<a name="table15465191317123"></a>
17<table><thead align="left"><tr id="row184651013131217"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p24653132122"><a name="p24653132122"></a><a name="p24653132122"></a>输入</p>16<table><thead align="left"><tr id="row184651013131217"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p24653132122"><a name="p24653132122"></a><a name="p24653132122"></a>输入</p>
18</th>17</th>
19<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p13465111311213"><a name="p13465111311213"></a><a name="p13465111311213"></a>Shape</p>18<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p13465111311213"><a name="p13465111311213"></a><a name="p13465111311213"></a>Shape</p>
@@ -18,9 +18,8 @@
18 18 
19本案例的算子规格如下:19本案例的算子规格如下:
20 20 
21-**表1** 算子规格21+**表1** 算子规格<a name="table101751125175213"></a>
22 22 
23-<a name="table101751125175213"></a>
24<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>23<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>
25</th>24</th>
26<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>25<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>
@@ -15,9 +15,8 @@
15 15 
16本案例的算子规格如下:16本案例的算子规格如下:
17 17 
18-**表1** 算子规格18+**表1** 算子规格<a name="table101751125175213"></a>
19 19 
20-<a name="table101751125175213"></a>
21<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>20<table><thead align="left"><tr id="row8175525185219"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p1417582516529"><a name="p1417582516529"></a><a name="p1417582516529"></a>输入</p>
22</th>21</th>
23<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>22<th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p417532575212"><a name="p417532575212"></a><a name="p417532575212"></a>Shape</p>
@@ -29,9 +29,8 @@
29- 算子类型(OpType):MATMUL\_LEAKYRELU29- 算子类型(OpType):MATMUL\_LEAKYRELU
30- 算子输入输出:30- 算子输入输出:
31 31 
32- **表1** MATMUL\_LEAKYRELU算子输入输出规格32+ **表1** MATMUL\_LEAKYRELU算子输入输出规格<a name="table4934296305"></a>
33 33 
34- <a name="table4934296305"></a>
35 <table><thead align="left"><tr id="row59358913304"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p5503181819300"><a name="p5503181819300"></a><a name="p5503181819300"></a><strong id="b1850331853010"><a name="b1850331853010"></a><a name="b1850331853010"></a>name</strong></p>34 <table><thead align="left"><tr id="row59358913304"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p5503181819300"><a name="p5503181819300"></a><a name="p5503181819300"></a><strong id="b1850331853010"><a name="b1850331853010"></a><a name="b1850331853010"></a>name</strong></p>
36 </th>35 </th>
37 <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1550381833017"><a name="p1550381833017"></a><a name="p1550381833017"></a><strong id="b7503171811309"><a name="b7503171811309"></a><a name="b7503171811309"></a>shape</strong></p>36 <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1550381833017"><a name="p1550381833017"></a><a name="p1550381833017"></a><strong id="b7503171811309"><a name="b7503171811309"></a><a name="b7503171811309"></a>shape</strong></p>
@@ -33,9 +33,8 @@
33 - 算子涉及Matmul左右矩阵在外部存储和内部存储间的数据搬运,查看Ascend C API参考中的数据搬运接口,需要使用[DataCopy](../../../../../api/SIMD-API/basic_api/data_move_guide/overview/data_move_concept.md)来实现数据搬运。33 - 算子涉及Matmul左右矩阵在外部存储和内部存储间的数据搬运,查看Ascend C API参考中的数据搬运接口,需要使用[DataCopy](../../../../../api/SIMD-API/basic_api/data_move_guide/overview/data_move_concept.md)来实现数据搬运。
34 - 计算过程涉及矩阵计算操作,查看Ascend C API参考中的矩阵计算相关接口,需要使用[Matmul高阶API](../../../../../api/SIMD-API/adv_api/cube_compute/Matmul_Kernel/Matmul_Kernel.md)实现矩阵乘计算。34 - 计算过程涉及矩阵计算操作,查看Ascend C API参考中的矩阵计算相关接口,需要使用[Matmul高阶API](../../../../../api/SIMD-API/adv_api/cube_compute/Matmul_Kernel/Matmul_Kernel.md)实现矩阵乘计算。
35 35 
36-**表1** AllGatherMatmulCustom算子规格36+**表1** AllGatherMatmulCustom算子规格<a name="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_table164881913121819"></a>
37 37 
38-<a name="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_table164881913121819"></a>
39<table><tbody><tr id="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_row1848911314188"><th class="firstcol" valign="top" id="mcps1.2.6.1.1"><p id="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"><a name="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"></a><a name="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"></a>算子类型(OpType)</p>38<table><tbody><tr id="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_row1848911314188"><th class="firstcol" valign="top" id="mcps1.2.6.1.1"><p id="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"><a name="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"></a><a name="zh-cn_topic_0000002400208581_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"></a>算子类型(OpType)</p>
40</th>39</th>
41<td class="cellrowborder" colspan="4" valign="top" headers="mcps1.2.6.1.1 "><p id="zh-cn_topic_0000002400208581_p13423114244112"><a name="zh-cn_topic_0000002400208581_p13423114244112"></a><a name="zh-cn_topic_0000002400208581_p13423114244112"></a><span>AllGatherMatmulCustom</span></p>40<td class="cellrowborder" colspan="4" valign="top" headers="mcps1.2.6.1.1 "><p id="zh-cn_topic_0000002400208581_p13423114244112"><a name="zh-cn_topic_0000002400208581_p13423114244112"></a><a name="zh-cn_topic_0000002400208581_p13423114244112"></a><span>AllGatherMatmulCustom</span></p>
@@ -74,9 +74,8 @@
74 74 
75## 参数说明<a name="zh-cn_topic_0000002298767897_section139621338103518"></a>75## 参数说明<a name="zh-cn_topic_0000002298767897_section139621338103518"></a>
76 76 
77-**表1** SparseMatmulType类型参数说明77+**表1** SparseMatmulType类型参数说明<a name="zh-cn_topic_0000002298767897_table3658657131020"></a>
78 78 
79-<a name="zh-cn_topic_0000002298767897_table3658657131020"></a>
80<table><thead align="left"><tr id="zh-cn_topic_0000002298767897_row8658257191010"><th class="cellrowborder" valign="top" width="21.21%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002298767897_p191082218111"><a name="zh-cn_topic_0000002298767897_p191082218111"></a><a name="zh-cn_topic_0000002298767897_p191082218111"></a>参数</p>79<table><thead align="left"><tr id="zh-cn_topic_0000002298767897_row8658257191010"><th class="cellrowborder" valign="top" width="21.21%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002298767897_p191082218111"><a name="zh-cn_topic_0000002298767897_p191082218111"></a><a name="zh-cn_topic_0000002298767897_p191082218111"></a>参数</p>
81</th>80</th>
82<th class="cellrowborder" valign="top" width="78.79%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002298767897_p1554812516113"><a name="zh-cn_topic_0000002298767897_p1554812516113"></a><a name="zh-cn_topic_0000002298767897_p1554812516113"></a>说明</p>81<th class="cellrowborder" valign="top" width="78.79%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002298767897_p1554812516113"><a name="zh-cn_topic_0000002298767897_p1554812516113"></a><a name="zh-cn_topic_0000002298767897_p1554812516113"></a>说明</p>
@@ -2,9 +2,8 @@
2 2 
3除了前述[基础知识](../basics.md)和[算子实现](../operator_impl.md)中介绍的基本计算能力外,Matmul矩阵编程还提供了适用于不同场景的处理能力及多种功能,具体场景和功能列于下表中,详细内容请见后续章节。3除了前述[基础知识](../basics.md)和[算子实现](../operator_impl.md)中介绍的基本计算能力外,Matmul矩阵编程还提供了适用于不同场景的处理能力及多种功能,具体场景和功能列于下表中,详细内容请见后续章节。
4 4 
5-**表1** Matmul功能特性表5+**表1** Matmul功能特性表<a name="zh-cn_topic_0000002298078557_table1220612410208"></a>
6 6 
7-<a name="zh-cn_topic_0000002298078557_table1220612410208"></a>
8<table><thead align="left"><tr id="zh-cn_topic_0000002298078557_row19206141132015"><th class="cellrowborder" valign="top" width="23.64%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002298078557_p52061441122017"><a name="zh-cn_topic_0000002298078557_p52061441122017"></a><a name="zh-cn_topic_0000002298078557_p52061441122017"></a>特性描述</p>7<table><thead align="left"><tr id="zh-cn_topic_0000002298078557_row19206141132015"><th class="cellrowborder" valign="top" width="23.64%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002298078557_p52061441122017"><a name="zh-cn_topic_0000002298078557_p52061441122017"></a><a name="zh-cn_topic_0000002298078557_p52061441122017"></a>特性描述</p>
9</th>8</th>
10<th class="cellrowborder" valign="top" width="76.36%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002298078557_p1206124122015"><a name="zh-cn_topic_0000002298078557_p1206124122015"></a><a name="zh-cn_topic_0000002298078557_p1206124122015"></a>功能简介</p>9<th class="cellrowborder" valign="top" width="76.36%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002298078557_p1206124122015"><a name="zh-cn_topic_0000002298078557_p1206124122015"></a><a name="zh-cn_topic_0000002298078557_p1206124122015"></a>功能简介</p>
@@ -84,9 +83,8 @@
84</tbody>83</tbody>
85</table>84</table>
86 85 
87-**表2** BatchMatmu功能l特性表86+**表2** BatchMatmul功能特性表<a name="zh-cn_topic_0000002298078557_table1558122415575"></a>
88 87 
89-<a name="zh-cn_topic_0000002298078557_table1558122415575"></a>
90<table><thead align="left"><tr id="zh-cn_topic_0000002298078557_row1255852419575"><th class="cellrowborder" valign="top" width="24.04%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002298078557_p4558162445718"><a name="zh-cn_topic_0000002298078557_p4558162445718"></a><a name="zh-cn_topic_0000002298078557_p4558162445718"></a>特性描述</p>88<table><thead align="left"><tr id="zh-cn_topic_0000002298078557_row1255852419575"><th class="cellrowborder" valign="top" width="24.04%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002298078557_p4558162445718"><a name="zh-cn_topic_0000002298078557_p4558162445718"></a><a name="zh-cn_topic_0000002298078557_p4558162445718"></a>特性描述</p>
91</th>89</th>
92<th class="cellrowborder" valign="top" width="75.96000000000001%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002298078557_p105581124185715"><a name="zh-cn_topic_0000002298078557_p105581124185715"></a><a name="zh-cn_topic_0000002298078557_p105581124185715"></a>功能简介</p>90<th class="cellrowborder" valign="top" width="75.96000000000001%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002298078557_p105581124185715"><a name="zh-cn_topic_0000002298078557_p105581124185715"></a><a name="zh-cn_topic_0000002298078557_p105581124185715"></a>功能简介</p>
@@ -30,9 +30,7 @@ Matmul量化/反量化包含两种模式:同一系数的量化/反量化模式
30 30 
31需要对矩阵计算结果进行量化/反量化操作的场景,当前该场景下,Matmul输入输出矩阵支持的数据类型如下表所示。31需要对矩阵计算结果进行量化/反量化操作的场景,当前该场景下,Matmul输入输出矩阵支持的数据类型如下表所示。
32 32 
33-**表2** Matmul量化/反量化支持的数据类型33+**表2** Matmul量化/反量化支持的数据类型<a name="zh-cn_topic_0000002298767893_table1996113269499"></a>
34- 
35-<a name="zh-cn_topic_0000002298767893_table1996113269499"></a>
36 34 
37| A矩阵 | B矩阵 | C矩阵 | 支持平台 |35| A矩阵 | B矩阵 | C矩阵 | 支持平台 |
38| --- | --- | --- | --- |36| --- | --- | --- | --- |
@@ -19,9 +19,7 @@ MX格式是一种块数据格式,若干个数据可以组成一个块(或者
19 19 
20MX格式的数据类型包含多种,例如MXFP8、MXFP4、MXFP16、MXINT4等。下表列举了[MxMatmul场景](#zh-cn_topic_0000002270097206_section310824820358)(全称Microscaling Matmul)支持的数据类型。20MX格式的数据类型包含多种,例如MXFP8、MXFP4、MXFP16、MXINT4等。下表列举了[MxMatmul场景](#zh-cn_topic_0000002270097206_section310824820358)(全称Microscaling Matmul)支持的数据类型。
21 21 
22-**表1** MxMatmul支持MX格式的数据类型22+**表1** MxMatmul支持MX格式的数据类型<a name="zh-cn_topic_0000002270097206_table5383144710452"></a>
23- 
24-<a name="zh-cn_topic_0000002270097206_table5383144710452"></a>
25 23 
26| 数据类型 | 私有元素数据类型 | 私有元素位宽(d) | 块大小(k) | 共享缩放因子数据类型 | 共享缩放因子位宽(w) |24| 数据类型 | 私有元素数据类型 | 私有元素位宽(d) | 块大小(k) | 共享缩放因子数据类型 | 共享缩放因子位宽(w) |
27| --- | --- | --- | --- | --- | --- |25| --- | --- | --- | --- | --- | --- |
@@ -205,9 +203,7 @@ Host侧自动获取Tiling参数的关键步骤介绍如下:
205 203 
206## 参数说明204## 参数说明
207 205 
208-**表3** MatmulTypeWithScale参数说明206+**表3** MatmulTypeWithScale参数说明<a name="zh-cn_topic_0000002270097206_table14759942142014"></a>
209- 
210-<a name="zh-cn_topic_0000002270097206_table14759942142014"></a>
211 207 
212| 参数 | 说明 |208| 参数 | 说明 |
213| --- | --- |209| --- | --- |
@@ -7,8 +7,8 @@
7**图1** 矩阵编程流程示意图<a name="zh-cn_topic_0000001622514006_fig264116391245"></a> 7**图1** 矩阵编程流程示意图<a name="zh-cn_topic_0000001622514006_fig264116391245"></a>
8![](../../../figures/mat_prog_flow.png "矩阵编程流程示意图")8![](../../../figures/mat_prog_flow.png "矩阵编程流程示意图")
9 9 
10-host侧自动获取Tiling参数的关键步骤介绍如下:10+host侧自动获取Tiling参数的关键步骤介绍如下:<a name="zh-cn_topic_0000001622514006_li19113114819525"></a>
11-<a name="zh-cn_topic_0000001622514006_li19113114819525"></a>11+ 
121. **创建Tiling对象**。121. **创建Tiling对象**。
13 13 
14 ```14 ```
@@ -58,8 +58,7 @@ host侧自动获取Tiling参数的关键步骤介绍如下:
58 tilingData.SaveToBuffer(tilingBuf, tcubeTilingSize);58 tilingData.SaveToBuffer(tilingBuf, tcubeTilingSize);
59 ```59 ```
60 60 
61-<a name="zh-cn_topic_0000001622514006_li1032116474330"></a>61+核函数(Kernel)侧使用Matmul API矩阵乘运算的具体步骤如下:<a name="zh-cn_topic_0000001622514006_li1032116474330"></a>
62-核函数(Kernel)侧使用Matmul API矩阵乘运算的具体步骤如下:
63 62 
641. **创建Matmul对象。**631. **创建Matmul对象。**
65 64 
@@ -170,9 +170,8 @@ extern "C" __global__ __aicore__ void matmul_custom(__gm__ uint8_t* a, __gm__ ui
170**图4** Matmul算子的计算流程图<a name="zh-cn_topic_0000002135641293_fig5645174474016"></a> 170**图4** Matmul算子的计算流程图<a name="zh-cn_topic_0000002135641293_fig5645174474016"></a>
171![](../../../figures/mm_flow.png "Matmul算子的计算流程图")171![](../../../figures/mm_flow.png "Matmul算子的计算流程图")
172 172 
173-**表1** Ascend C Matmul算子设计规格173+**表1** Ascend C Matmul算子设计规格<a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_table164881913121819"></a>
174 174 
175-<a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_table164881913121819"></a>
176<table><tbody><tr id="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_row1848911314188"><th class="firstcol" valign="top" id="mcps1.2.6.1.1"><p id="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"><a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"></a><a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"></a>算子类型(OpType)</p>175<table><tbody><tr id="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_row1848911314188"><th class="firstcol" valign="top" id="mcps1.2.6.1.1"><p id="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"><a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"></a><a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p54891613141818"></a>算子类型(OpType)</p>
177</th>176</th>
178<td class="cellrowborder" colspan="4" valign="top" headers="mcps1.2.6.1.1 "><p id="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p1748971311817"><a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p1748971311817"></a><a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p1748971311817"></a>Matmul</p>177<td class="cellrowborder" colspan="4" valign="top" headers="mcps1.2.6.1.1 "><p id="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p1748971311817"><a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p1748971311817"></a><a name="zh-cn_topic_0000002135641293_zh-cn_topic_0000001514387805_zh-cn_topic_0229825559_zh-cn_topic_0229823837_zh-cn_topic_0211294710_p1748971311817"></a>Matmul</p>
@@ -48,9 +48,8 @@
48- 算子类型(OpType):Add48- 算子类型(OpType):Add
49- 算子输入输出:49- 算子输入输出:
50 50 
51- **表1** Add算子输入输出规格51+ **表1** Add算子输入输出规格<a name="table4934296305"></a>
52 52 
53- <a name="table4934296305"></a>
54 <table><thead align="left"><tr id="row59358913304"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p5503181819300"><a name="p5503181819300"></a><a name="p5503181819300"></a><strong id="b1850331853010"><a name="b1850331853010"></a><a name="b1850331853010"></a>name</strong></p>53 <table><thead align="left"><tr id="row59358913304"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p5503181819300"><a name="p5503181819300"></a><a name="p5503181819300"></a><strong id="b1850331853010"><a name="b1850331853010"></a><a name="b1850331853010"></a>name</strong></p>
55 </th>54 </th>
56 <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1550381833017"><a name="p1550381833017"></a><a name="p1550381833017"></a><strong id="b7503171811309"><a name="b7503171811309"></a><a name="b7503171811309"></a>shape</strong></p>55 <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1550381833017"><a name="p1550381833017"></a><a name="p1550381833017"></a><strong id="b7503171811309"><a name="b7503171811309"></a><a name="b7503171811309"></a>shape</strong></p>
@@ -215,8 +214,7 @@ __aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t
215 214 
216基于矢量编程范式,将核函数(Kernel)的实现分为3个基本任务:CopyIn,Compute,CopyOut。再根据编程范式上面的算法分析,将整个计算拆分成三个Stage,用户单独编写每个Stage的代码,三阶段流程示意图参见[图3](#zh-cn_topic_0000002201157438_fig4134406304),具体流程如下:215基于矢量编程范式,将核函数(Kernel)的实现分为3个基本任务:CopyIn,Compute,CopyOut。再根据编程范式上面的算法分析,将整个计算拆分成三个Stage,用户单独编写每个Stage的代码,三阶段流程示意图参见[图3](#zh-cn_topic_0000002201157438_fig4134406304),具体流程如下:
217 216 
218-1. Stage1:CopyIn实现。217+1. Stage1:CopyIn实现。<a id="copyin-implementation"></a>
219-<a id="copyin-implementation"></a>
220 218 
221 1. 使用[DataCopy](../../../../api/SIMD-API/basic_api/data_move_guide/overview/data_move_concept.md)接口将GlobalTensor数据拷贝到LocalTensor。219 1. 使用[DataCopy](../../../../api/SIMD-API/basic_api/data_move_guide/overview/data_move_concept.md)接口将GlobalTensor数据拷贝到LocalTensor。
222 2. 使用[EnQue](../../../../api/SIMD-API/basic_api/resource_management/TQue/EnQue.md)将LocalTensor放入UB(VECIN)的Queue中。220 2. 使用[EnQue](../../../../api/SIMD-API/basic_api/resource_management/TQue/EnQue.md)将LocalTensor放入UB(VECIN)的Queue中。
@@ -254,8 +252,7 @@ __aicore__ inline void Init(__gm__ uint8_t* x, __gm__ uint8_t* y, __gm__ uint8_t
254 inQueueY.FreeTensor(yLocal);252 inQueueY.FreeTensor(yLocal);
255 ```253 ```
256 254 
257-3. Stage3:CopyOut实现。255+3. Stage3:CopyOut实现。<a id="copyout-implementation"></a>
258-<a id="copyout-implementation"></a>
259 256 
260 1. 使用[DeQue](../../../../api/SIMD-API/basic_api/resource_management/TQue/DeQue.md)接口从UB(VECOUT)的Queue中取出LocalTensor。257 1. 使用[DeQue](../../../../api/SIMD-API/basic_api/resource_management/TQue/DeQue.md)接口从UB(VECOUT)的Queue中取出LocalTensor。
261 2. 使用[DataCopy](../../../../api/SIMD-API/basic_api/data_move_guide/overview/data_move_concept.md)接口将LocalTensor拷贝到GlobalTensor上。258 2. 使用[DataCopy](../../../../api/SIMD-API/basic_api/data_move_guide/overview/data_move_concept.md)接口将LocalTensor拷贝到GlobalTensor上。
@@ -9,9 +9,8 @@
9- 算子类型(OpType):Add9- 算子类型(OpType):Add
10- 算子输入输出:10- 算子输入输出:
11 11 
12- **表1** Add算子输入输出规格12+ **表1** Add算子输入输出规格<a name="table4934296305"></a>
13 13 
14- <a name="table4934296305"></a>
15 <table><thead align="left"><tr id="row59358913304"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p5503181819300"><a name="p5503181819300"></a><a name="p5503181819300"></a><strong id="b1850331853010"><a name="b1850331853010"></a><a name="b1850331853010"></a>name</strong></p>14 <table><thead align="left"><tr id="row59358913304"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p5503181819300"><a name="p5503181819300"></a><a name="p5503181819300"></a><strong id="b1850331853010"><a name="b1850331853010"></a><a name="b1850331853010"></a>name</strong></p>
16 </th>15 </th>
17 <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1550381833017"><a name="p1550381833017"></a><a name="p1550381833017"></a><strong id="b7503171811309"><a name="b7503171811309"></a><a name="b7503171811309"></a>shape</strong></p>16 <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1550381833017"><a name="p1550381833017"></a><a name="p1550381833017"></a><strong id="b7503171811309"><a name="b7503171811309"></a><a name="b7503171811309"></a>shape</strong></p>
@@ -30,9 +30,8 @@
30 30 
31为了体现同地址冲突的影响,上述场景关闭L2Cache hint。输入每行数据大小为1024字节(512个half),Tile为 \([128, 64]\),通过`DataCopyPad`完成GM到UB搬运。每个核都会完整载入一遍输入矩阵,N方向按64列切分,单次搬运数据量为`128 * 64 * 2`字节。原始实现中,所有核按相同的mBlockIdx顺序访问同一输入矩阵,容易在同一时刻访问相同的GM地址段;优化实现中,每个核按 \((mBlockIdx + blockIdx) \% numBlocks\)在块组内轮转访问顺序,降低同地址访问冲突概率。31为了体现同地址冲突的影响,上述场景关闭L2Cache hint。输入每行数据大小为1024字节(512个half),Tile为 \([128, 64]\),通过`DataCopyPad`完成GM到UB搬运。每个核都会完整载入一遍输入矩阵,N方向按64列切分,单次搬运数据量为`128 * 64 * 2`字节。原始实现中,所有核按相同的mBlockIdx顺序访问同一输入矩阵,容易在同一时刻访问相同的GM地址段;优化实现中,每个核按 \((mBlockIdx + blockIdx) \% numBlocks\)在块组内轮转访问顺序,降低同地址访问冲突概率。
32 32 
33-下文示意图中,M0-M47表示输入矩阵沿M方向切分后的48个分块,T0-T47表示时间轴上的48个搬运时刻。33+下文示意图中,M0-M47表示输入矩阵沿M方向切分后的48个分块,T0-T47表示时间轴上的48个搬运时刻。<a name="table69111151111819"></a>
34 34 
35-<a name="table69111151111819"></a>
36<table><thead align="left"><tr id="row199098514189"><th class="cellrowborder" valign="top" width="8.303319380901392%" id="mcps1.1.4.1.1"><p id="p7909145111187"><a name="p7909145111187"></a><a name="p7909145111187"></a>实现方案</p>35<table><thead align="left"><tr id="row199098514189"><th class="cellrowborder" valign="top" width="8.303319380901392%" id="mcps1.1.4.1.1"><p id="p7909145111187"><a name="p7909145111187"></a><a name="p7909145111187"></a>实现方案</p>
37</th>36</th>
38<th class="cellrowborder" valign="top" width="45.28375352866739%" id="mcps1.1.4.1.2"><p id="p18909175121819"><a name="p18909175121819"></a><a name="p18909175121819"></a>原始实现</p>37<th class="cellrowborder" valign="top" width="45.28375352866739%" id="mcps1.1.4.1.2"><p id="p18909175121819"><a name="p18909175121819"></a><a name="p18909175121819"></a>原始实现</p>
@@ -37,9 +37,8 @@ bank冲突主要可以分为以下三种场景:
37 37 
38 Vector指令的源操作数src和目的操作数dst同时读写到同一个bank时造成读写冲突,具体分析如下:38 Vector指令的源操作数src和目的操作数dst同时读写到同一个bank时造成读写冲突,具体分析如下:
39 39 
40- **表1** 读写冲突示例40+ **表1** 读写冲突示例<a name="table06741342154717"></a>
41 41 
42- <a name="table06741342154717"></a>
43 <table><thead align="left"><tr id="row1767434244710"><th class="cellrowborder" valign="top" width="5.9988002399520095%" id="mcps1.2.7.1.1"><p id="p17674144211470"><a name="p17674144211470"></a><a name="p17674144211470"></a>序号</p>42 <table><thead align="left"><tr id="row1767434244710"><th class="cellrowborder" valign="top" width="5.9988002399520095%" id="mcps1.2.7.1.1"><p id="p17674144211470"><a name="p17674144211470"></a><a name="p17674144211470"></a>序号</p>
44 </th>43 </th>
45 <th class="cellrowborder" valign="top" width="9.948010397920417%" id="mcps1.2.7.1.2"><p id="p9674114213473"><a name="p9674114213473"></a><a name="p9674114213473"></a>src地址</p>44 <th class="cellrowborder" valign="top" width="9.948010397920417%" id="mcps1.2.7.1.2"><p id="p9674114213473"><a name="p9674114213473"></a><a name="p9674114213473"></a>src地址</p>
@@ -87,9 +86,8 @@ bank冲突主要可以分为以下三种场景:
87 86 
88 Vector指令目的操作数dst对应的8个DataBlock(block0-block7)同时写到一个bank group时造成写写冲突,具体分析如下:87 Vector指令目的操作数dst对应的8个DataBlock(block0-block7)同时写到一个bank group时造成写写冲突,具体分析如下:
89 88 
90- **表2** 写写冲突示例89+ **表2** 写写冲突示例<a name="table11121346153814"></a>
91 90 
92- <a name="table11121346153814"></a>
93 <table><thead align="left"><tr id="row212044623812"><th class="cellrowborder" valign="top" width="5.789726356216995%" id="mcps1.2.9.1.1"><p id="p13120184619386"><a name="p13120184619386"></a><a name="p13120184619386"></a>序号</p>91 <table><thead align="left"><tr id="row212044623812"><th class="cellrowborder" valign="top" width="5.789726356216995%" id="mcps1.2.9.1.1"><p id="p13120184619386"><a name="p13120184619386"></a><a name="p13120184619386"></a>序号</p>
94 </th>92 </th>
95 <th class="cellrowborder" valign="top" width="7.959673547767644%" id="mcps1.2.9.1.2"><p id="p2012004612384"><a name="p2012004612384"></a><a name="p2012004612384"></a>dst地址</p>93 <th class="cellrowborder" valign="top" width="7.959673547767644%" id="mcps1.2.9.1.2"><p id="p2012004612384"><a name="p2012004612384"></a><a name="p2012004612384"></a>dst地址</p>
@@ -148,9 +146,8 @@ bank冲突主要可以分为以下三种场景:
148- 读读冲突146- 读读冲突
149 - Vector指令多个源操作数同时读到同一个bank group时造成读读冲突,具体分析如下:147 - Vector指令多个源操作数同时读到同一个bank group时造成读读冲突,具体分析如下:
150 148 
151- **表3** 双src场景读读冲突示例149+ **表3** 双src场景读读冲突示例<a name="table1318881512255"></a>
152 150 
153- <a name="table1318881512255"></a>
154 <table><thead align="left"><tr id="row8188715102517"><th class="cellrowborder" valign="top" width="6.5786842631473705%" id="mcps1.2.7.1.1"><p id="p111881154254"><a name="p111881154254"></a><a name="p111881154254"></a>序号</p>151 <table><thead align="left"><tr id="row8188715102517"><th class="cellrowborder" valign="top" width="6.5786842631473705%" id="mcps1.2.7.1.1"><p id="p111881154254"><a name="p111881154254"></a><a name="p111881154254"></a>序号</p>
155 </th>152 </th>
156 <th class="cellrowborder" valign="top" width="12.787442511497702%" id="mcps1.2.7.1.2"><p id="p14188915172513"><a name="p14188915172513"></a><a name="p14188915172513"></a>src0地址</p>153 <th class="cellrowborder" valign="top" width="12.787442511497702%" id="mcps1.2.7.1.2"><p id="p14188915172513"><a name="p14188915172513"></a><a name="p14188915172513"></a>src0地址</p>
@@ -196,9 +193,8 @@ bank冲突主要可以分为以下三种场景:
196 193 
197 - Vector指令某一个源操作数对应的8个DataBlock(block0-block7)读到同一个bank group时造成读读冲突,具体分析如下:194 - Vector指令某一个源操作数对应的8个DataBlock(block0-block7)读到同一个bank group时造成读读冲突,具体分析如下:
198 195 
199- **表4** 单src场景读读冲突示例196+ **表4** 单src场景读读冲突示例<a name="table332972534717"></a>
200 197 
201- <a name="table332972534717"></a>
202 <table><thead align="left"><tr id="row832822516476"><th class="cellrowborder" valign="top" width="6.117919521374119%" id="mcps1.2.9.1.1"><p id="p157111614175011"><a name="p157111614175011"></a><a name="p157111614175011"></a>序号</p>198 <table><thead align="left"><tr id="row832822516476"><th class="cellrowborder" valign="top" width="6.117919521374119%" id="mcps1.2.9.1.1"><p id="p157111614175011"><a name="p157111614175011"></a><a name="p157111614175011"></a>序号</p>
203 </th>199 </th>
204 <th class="cellrowborder" valign="top" width="11.560358969410403%" id="mcps1.2.9.1.2"><p id="p33287255478"><a name="p33287255478"></a><a name="p33287255478"></a>src地址</p>200 <th class="cellrowborder" valign="top" width="11.560358969410403%" id="mcps1.2.9.1.2"><p id="p33287255478"><a name="p33287255478"></a><a name="p33287255478"></a>src地址</p>
@@ -264,9 +260,8 @@ bank冲突主要可以分为以下三种场景:
264 260 
265- **优化计算逻辑**261- **优化计算逻辑**
266 262 
267- 对一个shape为\(8, 16, 16\)的输入做\(1, 0, 2\)的transpose操作,输出shape为\(16, 8, 16\)。通过将计算逻辑由“跳读,连续写”修改为“连续读,跳写”可避免冲突问题。实现方案对比如下:263+ 对一个shape为\(8, 16, 16\)的输入做\(1, 0, 2\)的transpose操作,输出shape为\(16, 8, 16\)。通过将计算逻辑由“跳读,连续写”修改为“连续读,跳写”可避免冲突问题。实现方案对比如下:<a name="table12921549195512"></a>
268 264 
269- <a name="table12921549195512"></a>
270 <table><thead align="left"><tr id="row1229364945511"><th class="cellrowborder" valign="top" width="6.813978389954251%" id="mcps1.1.4.1.1"><p id="p2081249145715"><a name="p2081249145715"></a><a name="p2081249145715"></a>实现方案</p>265 <table><thead align="left"><tr id="row1229364945511"><th class="cellrowborder" valign="top" width="6.813978389954251%" id="mcps1.1.4.1.1"><p id="p2081249145715"><a name="p2081249145715"></a><a name="p2081249145715"></a>实现方案</p>
271 </th>266 </th>
272 <th class="cellrowborder" valign="top" width="42.6652389759564%" id="mcps1.1.4.1.2"><p id="p2029374985519"><a name="p2029374985519"></a><a name="p2029374985519"></a>原始实现</p>267 <th class="cellrowborder" valign="top" width="42.6652389759564%" id="mcps1.1.4.1.2"><p id="p2029374985519"><a name="p2029374985519"></a><a name="p2029374985519"></a>原始实现</p>
@@ -50,9 +50,8 @@ UB_ADDR[18:0] = {BANK_DEPTH[8:0], BANK[0], BG[2:0], BANK_WIDTH[4:0]}
50 50 
51 同时读写到同一个bank时造成读写冲突,具体分析如下:51 同时读写到同一个bank时造成读写冲突,具体分析如下:
52 52 
53- **表1** 一读一写冲突示例53+ **表1** 一读一写冲突示例<a name="table1789735214010"></a>
54 54 
55- <a name="table1789735214010"></a>
56 <table><thead align="left"><tr id="row20897752164015"><th class="cellrowborder" valign="top" width="5.9988002399520095%" id="mcps1.2.7.1.1"><p id="p13897195244019"><a name="p13897195244019"></a><a name="p13897195244019"></a>序号</p>55 <table><thead align="left"><tr id="row20897752164015"><th class="cellrowborder" valign="top" width="5.9988002399520095%" id="mcps1.2.7.1.1"><p id="p13897195244019"><a name="p13897195244019"></a><a name="p13897195244019"></a>序号</p>
57 </th>56 </th>
58 <th class="cellrowborder" valign="top" width="9.948010397920417%" id="mcps1.2.7.1.2"><p id="p3897155224010"><a name="p3897155224010"></a><a name="p3897155224010"></a>src地址</p>57 <th class="cellrowborder" valign="top" width="9.948010397920417%" id="mcps1.2.7.1.2"><p id="p3897155224010"><a name="p3897155224010"></a><a name="p3897155224010"></a>src地址</p>
@@ -263,9 +262,8 @@ UB_ADDR[18:0] = {BANK_DEPTH[8:0], BANK[0], BG[2:0], BANK_WIDTH[4:0]}
263 262 
264- 优化计算逻辑263- 优化计算逻辑
265 264 
266- 对一个数据类型为float,shape为\(8, 64\)的输入每个元素加1。通过将计算逻辑由逐列计算改为逐行计算可避免同一Repeat下的冲突问题,实现方案对比如下:265+ 对一个数据类型为float,shape为\(8, 64\)的输入每个元素加1。通过将计算逻辑由逐列计算改为逐行计算可避免同一Repeat下的冲突问题,实现方案对比如下:<a name="table12921549195512"></a>
267 266 
268- <a name="table12921549195512"></a>
269 <table><thead align="left"><tr id="row1229364945511"><th class="cellrowborder" valign="top" width="6.813978389954251%" id="mcps1.1.4.1.1"><p id="p2081249145715"><a name="p2081249145715"></a><a name="p2081249145715"></a>实现方案</p>267 <table><thead align="left"><tr id="row1229364945511"><th class="cellrowborder" valign="top" width="6.813978389954251%" id="mcps1.1.4.1.1"><p id="p2081249145715"><a name="p2081249145715"></a><a name="p2081249145715"></a>实现方案</p>
270 </th>268 </th>
271 <th class="cellrowborder" valign="top" width="42.6652389759564%" id="mcps1.1.4.1.2"><p id="p2029374985519"><a name="p2029374985519"></a><a name="p2029374985519"></a>原始实现</p>269 <th class="cellrowborder" valign="top" width="42.6652389759564%" id="mcps1.1.4.1.2"><p id="p2029374985519"><a name="p2029374985519"></a><a name="p2029374985519"></a>原始实现</p>
@@ -38,9 +38,8 @@ L2 Cache的带宽相比GM的带宽有数倍的提升,因此当数据命中L2 C
38 38 
39Ascend C基于GlobalTensor提供了SetL2CacheHint接口,用户可以根据需要指定CacheMode。39Ascend C基于GlobalTensor提供了SetL2CacheHint接口,用户可以根据需要指定CacheMode。
40 40 
41-考虑如下场景,构造两个Tensor的相加计算,输入x、y和输出z的Shape均为\(8192, 8192\)。下文完整样例请参考[设置合理L2 CacheMode样例](../../../../../../examples/01_simd_cpp_api/05_best_practices/00_vector_compute/add_high_performance)。41+考虑如下场景,构造两个Tensor的相加计算,输入x、y和输出z的Shape均为\(8192, 8192\)。下文完整样例请参考[设置合理L2 CacheMode样例](../../../../../../examples/01_simd_cpp_api/05_best_practices/00_vector_compute/add_high_performance)。<a name="table12921549195512"></a>
42 42 
43-<a name="table12921549195512"></a>
44<table><thead align="left"><tr id="row1229364945511"><th class="cellrowborder" valign="top" width="6.813978389954251%" id="mcps1.1.4.1.1"><p id="p2081249145715"><a name="p2081249145715"></a><a name="p2081249145715"></a>实现方案</p>43<table><thead align="left"><tr id="row1229364945511"><th class="cellrowborder" valign="top" width="6.813978389954251%" id="mcps1.1.4.1.1"><p id="p2081249145715"><a name="p2081249145715"></a><a name="p2081249145715"></a>实现方案</p>
45</th>44</th>
46<th class="cellrowborder" valign="top" width="42.67497323079919%" id="mcps1.1.4.1.2"><p id="p2029374985519"><a name="p2029374985519"></a><a name="p2029374985519"></a>原始实现</p>45<th class="cellrowborder" valign="top" width="42.67497323079919%" id="mcps1.1.4.1.2"><p id="p2029374985519"><a name="p2029374985519"></a><a name="p2029374985519"></a>原始实现</p>
@@ -1,8 +1,7 @@
1# 优化建议总览表<a name="ZH-CN_TOPIC_0000002351966029"></a>1# 优化建议总览表<a name="ZH-CN_TOPIC_0000002351966029"></a>
2 2 
3-**表1** 性能优化建议总览表3+**表1** 性能优化建议总览表<a name="table1318114421218"></a>
4 4 
5-<a name="table1318114421218"></a>
6<table><thead align="left"><tr id="row8182944131211"><th class="cellrowborder" valign="top" width="12.55%" id="mcps1.2.4.1.1"><p id="p918218445121"><a name="p918218445121"></a><a name="p918218445121"></a>分类</p>5<table><thead align="left"><tr id="row8182944131211"><th class="cellrowborder" valign="top" width="12.55%" id="mcps1.2.4.1.1"><p id="p918218445121"><a name="p918218445121"></a><a name="p918218445121"></a>分类</p>
7</th>6</th>
8<th class="cellrowborder" valign="top" width="42.89%" id="mcps1.2.4.1.2"><p id="p1165171215245"><a name="p1165171215245"></a><a name="p1165171215245"></a>分类描述</p>7<th class="cellrowborder" valign="top" width="42.89%" id="mcps1.2.4.1.2"><p id="p1165171215245"><a name="p1165171215245"></a><a name="p1165171215245"></a>分类描述</p>
@@ -26,9 +26,8 @@
26 26 
27【性能数据】27【性能数据】
28 28 
29-输入shape为30000,数据类型为float时,如下示例的性能数据对比如下,数据单位为cycle,使用GetSystemCycle接口获取。29+输入shape为30000,数据类型为float时,如下示例的性能数据对比如下,数据单位为cycle,使用GetSystemCycle接口获取。<a name="table1240852315316"></a>
30 30 
31-<a name="table1240852315316"></a>
32<table><thead align="left"><tr id="row140811238537"><th class="cellrowborder" valign="top" width="50%" id="mcps1.1.3.1.1"><p id="p12319202885319"><a name="p12319202885319"></a><a name="p12319202885319"></a>二分累加方案</p>31<table><thead align="left"><tr id="row140811238537"><th class="cellrowborder" valign="top" width="50%" id="mcps1.1.3.1.1"><p id="p12319202885319"><a name="p12319202885319"></a><a name="p12319202885319"></a>二分累加方案</p>
33</th>32</th>
34<th class="cellrowborder" valign="top" width="50%" id="mcps1.1.3.1.2"><p id="p1788511331538"><a name="p1788511331538"></a><a name="p1788511331538"></a>ReduceRepeat单指令操作</p>33<th class="cellrowborder" valign="top" width="50%" id="mcps1.1.3.1.2"><p id="p1788511331538"><a name="p1788511331538"></a><a name="p1788511331538"></a>ReduceRepeat单指令操作</p>
@@ -120,9 +119,8 @@ __aicore__ inline void ReduceRepeatSumImpl(const AscendC::LocalTensor<float>& ds
120 119 
121输入shape为256,数据类型为float。如下示例的性能数据如下:120输入shape为256,数据类型为float。如下示例的性能数据如下:
122 121 
123-**表1** 两次ReduceRepeat、三次ReduceDataBlock、一次ReduceDataBlock加一次ReduceRepeat,三种归约操作的性能数据(循环100次的时间总和)122+**表1** 两次ReduceRepeat、三次ReduceDataBlock、一次ReduceDataBlock加一次ReduceRepeat,三种归约操作的性能数据(循环100次的时间总和)<a name="table174575114592"></a>
124 123 
125-<a name="table174575114592"></a>
126<table><thead align="left"><tr id="row3758185125915"><th class="cellrowborder" valign="top" width="33.33333333333333%" id="mcps1.2.4.1.1"><p id="p57581151175915"><a name="p57581151175915"></a><a name="p57581151175915"></a>两次ReduceRepeat</p>124<table><thead align="left"><tr id="row3758185125915"><th class="cellrowborder" valign="top" width="33.33333333333333%" id="mcps1.2.4.1.1"><p id="p57581151175915"><a name="p57581151175915"></a><a name="p57581151175915"></a>两次ReduceRepeat</p>
127</th>125</th>
128<th class="cellrowborder" valign="top" width="33.33333333333333%" id="mcps1.2.4.1.2"><p id="p1960391014012"><a name="p1960391014012"></a><a name="p1960391014012"></a>三次ReduceDataBlock</p>126<th class="cellrowborder" valign="top" width="33.33333333333333%" id="mcps1.2.4.1.2"><p id="p1960391014012"><a name="p1960391014012"></a><a name="p1960391014012"></a>三次ReduceDataBlock</p>
@@ -30,9 +30,8 @@
30 30 
313. 明确函数名称和参数。313. 明确函数名称和参数。
32 - 自定义核函数(Kernel)名称,本样例中核函数(Kernel)命名为gather\_custom。32 - 自定义核函数(Kernel)名称,本样例中核函数(Kernel)命名为gather\_custom。
33- - 通过分析算子的输入和输出,使用模板参数来支持不同的输入输出数据类型。33+ - 通过分析算子的输入和输出,使用模板参数来支持不同的输入输出数据类型。<a name="table4135123919714"></a>
34 34 
35- <a name="table4135123919714"></a>
36 <table><thead align="left"><tr id="row6135439176"><th class="cellrowborder" valign="top" width="31.619999999999997%" id="mcps1.1.4.1.1"><p id="p213510395713"><a name="p213510395713"></a><a name="p213510395713"></a>模板参数名</p>35 <table><thead align="left"><tr id="row6135439176"><th class="cellrowborder" valign="top" width="31.619999999999997%" id="mcps1.1.4.1.1"><p id="p213510395713"><a name="p213510395713"></a><a name="p213510395713"></a>模板参数名</p>
37 </th>36 </th>
38 <th class="cellrowborder" valign="top" width="29.2%" id="mcps1.1.4.1.2"><p id="p531020292080"><a name="p531020292080"></a><a name="p531020292080"></a>模板参数类型</p>37 <th class="cellrowborder" valign="top" width="29.2%" id="mcps1.1.4.1.2"><p id="p531020292080"><a name="p531020292080"></a><a name="p531020292080"></a>模板参数类型</p>
@@ -58,9 +57,8 @@
58 </tbody>57 </tbody>
59 </table>58 </table>
60 59 
61- 函数入参定义如下:60+ 函数入参定义如下:<a name="table1353353310716"></a>
62 61 
63- <a name="table1353353310716"></a>
64 <table><thead align="left"><tr id="row125311633871"><th class="cellrowborder" valign="top" width="23.362336233623363%" id="mcps1.1.4.1.1"><p id="p115315333718"><a name="p115315333718"></a><a name="p115315333718"></a>参数名</p>62 <table><thead align="left"><tr id="row125311633871"><th class="cellrowborder" valign="top" width="23.362336233623363%" id="mcps1.1.4.1.1"><p id="p115315333718"><a name="p115315333718"></a><a name="p115315333718"></a>参数名</p>
65 </th>63 </th>
66 <th class="cellrowborder" valign="top" width="17.72177217721772%" id="mcps1.1.4.1.2"><p id="p135316331477"><a name="p135316331477"></a><a name="p135316331477"></a>参数类型</p>64 <th class="cellrowborder" valign="top" width="17.72177217721772%" id="mcps1.1.4.1.2"><p id="p135316331477"><a name="p135316331477"></a><a name="p135316331477"></a>参数类型</p>
@@ -116,9 +114,8 @@
116- 算子类型(OpType):Gather114- 算子类型(OpType):Gather
117- 算子输入输出:115- 算子输入输出:
118 116 
119- **表1** Gather算子输入输出规格117+ **表1** Gather算子输入输出规格<a name="table4934296305"></a>
120 118 
121- <a name="table4934296305"></a>
122 <table><thead align="left"><tr id="row59358913304"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p5503181819300"><a name="p5503181819300"></a><a name="p5503181819300"></a><strong id="b1850331853010"><a name="b1850331853010"></a><a name="b1850331853010"></a>name</strong></p>119 <table><thead align="left"><tr id="row59358913304"><th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.1"><p id="p5503181819300"><a name="p5503181819300"></a><a name="p5503181819300"></a><strong id="b1850331853010"><a name="b1850331853010"></a><a name="b1850331853010"></a>name</strong></p>
123 </th>120 </th>
124 <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1550381833017"><a name="p1550381833017"></a><a name="p1550381833017"></a><strong id="b7503171811309"><a name="b7503171811309"></a><a name="b7503171811309"></a>shape</strong></p>121 <th class="cellrowborder" valign="top" width="25%" id="mcps1.2.5.1.2"><p id="p1550381833017"><a name="p1550381833017"></a><a name="p1550381833017"></a><strong id="b7503171811309"><a name="b7503171811309"></a><a name="b7503171811309"></a>shape</strong></p>
@@ -147,9 +147,7 @@ cmake --build build_out --target binary package -j$(nproc)
147export ASCENDC_BUILD_LOG_DIR=/home/build_log/147export ASCENDC_BUILD_LOG_DIR=/home/build_log/
148```148```
149 149 
150-<a id="算子包部署"></a>150+## 算子包部署<a id="operator-package-deployment"></a><a id="算子包部署"></a>
151- 
152-## 算子包部署<a id="operator-package-deployment"></a>
153 151 
154编译生成的算子安装包需要部署到运行环境。152编译生成的算子安装包需要部署到运行环境。
155 153 
@@ -184,9 +182,7 @@ source /usr/local/Ascend/cann/set_env.sh
184> - 指定目录安装后,在调用算子前需要执行`source <path>/vendors/<vendor_name>/bin/set_env.bash`。该脚本会将算子包目录`<path>/vendors/<vendor_name>`添加到`ASCEND_CUSTOM_OPP_PATH`环境变量的首部,供框架在运行时查找自定义算子包;同时将`<path>/vendors/<vendor_name>/op_api/lib/`添加到`LD_LIBRARY_PATH`环境变量的首部,供系统动态链接器查找算子动态库。182> - 指定目录安装后,在调用算子前需要执行`source <path>/vendors/<vendor_name>/bin/set_env.bash`。该脚本会将算子包目录`<path>/vendors/<vendor_name>`添加到`ASCEND_CUSTOM_OPP_PATH`环境变量的首部,供框架在运行时查找自定义算子包;同时将`<path>/vendors/<vendor_name>/op_api/lib/`添加到`LD_LIBRARY_PATH`环境变量的首部,供系统动态链接器查找算子动态库。
185> - 算子包要求架构与系统架构一致,且满足glibc版本依赖,否则可使用`--force`强制安装。183> - 算子包要求架构与系统架构一致,且满足glibc版本依赖,否则可使用`--force`强制安装。
186 184 
187-<a id="部署后的目录结构"></a>185+### 部署后的目录结构<a id="部署后的目录结构"></a>
188- 
189-### 部署后的目录结构
190 186 
191以默认安装场景为例:187以默认安装场景为例:
192 188 
@@ -224,9 +220,7 @@ source /usr/local/Ascend/cann/set_env.sh
224│ └── version.info //版本信息220│ └── version.info //版本信息
225```221```
226 222 
227-<a id="配置自定义算子优先级"></a>223+### 配置自定义算子优先级<a id="配置自定义算子优先级"></a>
228- 
229-### 配置自定义算子优先级
230 224 
231多个算子包共存时,如果包含相同OpType的算子,优先级高的算子包中的版本生效。225多个算子包共存时,如果包含相同OpType的算子,优先级高的算子包中的版本生效。
232 226 
@@ -315,9 +309,7 @@ source /usr/local/Ascend/cann/set_env.sh
315 309 
316---310---
317 311 
318-<a id="编译组织"></a>312+## 编译组织(进阶)<a id="advanced-build-organization"></a><a id="编译组织"></a>
319- 
320-## 编译组织(进阶)<a id="advanced-build-organization"></a>
321 313 
322算子工程通过CMake组织编译。如需自行组织工程或用CANN提供的CMake modules定制编译流程,可参考以下内容。314算子工程通过CMake组织编译。如需自行组织工程或用CANN提供的CMake modules定制编译流程,可参考以下内容。
323 315 
@@ -616,9 +608,7 @@ npu_op_package_add(${package_name}
616 608 
617完成编译组织之后,可按照[基本流程](#基本流程)的三个步骤完成算子编译与部署。609完成编译组织之后,可按照[基本流程](#基本流程)的三个步骤完成算子编译与部署。
618 610 
619-<a id="自定义编译选项"></a>611+## 自定义编译选项(进阶)<a id="custom-compile-options"></a><a id="自定义编译选项"></a>
620- 
621-## 自定义编译选项(进阶)<a id="custom-compile-options"></a>
622 612 
623在CMakeLists.txt中,用`npu_op_kernel_options`为核函数(Kernel)侧代码添加编译选项:613在CMakeLists.txt中,用`npu_op_kernel_options`为核函数(Kernel)侧代码添加编译选项:
624 614 
@@ -150,9 +150,7 @@ ccache --show-stats -v
150对比两次`ccache --show-stats -v`结果,可看到在机器B上远程缓存命中率`Remote storage Hits`显著提升。150对比两次`ccache --show-stats -v`结果,可看到在机器B上远程缓存命中率`Remote storage Hits`显著提升。
151 151 
152 152 
153-<a id="选择性编译"></a>153+## 选择性编译<a id="选择性编译"></a>
154- 
155-## 选择性编译
156 154 
157选择性编译的核心思路是:开发调试阶段只编译当前真正需要验证的部分,减少无关核函数(Kernel)变体、无关芯片和无关产物带来的编译开销。常用策略如下:155选择性编译的核心思路是:开发调试阶段只编译当前真正需要验证的部分,减少无关核函数(Kernel)变体、无关芯片和无关产物带来的编译开销。常用策略如下:
158 156 
@@ -2,9 +2,7 @@
2 2 
3SIMD与SIMT混合编程涉及多种类型的函数,它们之间遵循严格的调用关系和层级约束。本章节将介绍整体函数类型和调用关系,再详细展开每种函数的定义、调用语法及约束。3SIMD与SIMT混合编程涉及多种类型的函数,它们之间遵循严格的调用关系和层级约束。本章节将介绍整体函数类型和调用关系,再详细展开每种函数的定义、调用语法及约束。
4 4 
5-SIMD与SIMT混合编程中涉及的函数类型如下表所示:5+SIMD与SIMT混合编程中涉及的函数类型如下表所示:<a name="zh-cn_topic_0000002571578013_table1818191292017"></a>
6- 
7-<a name="zh-cn_topic_0000002571578013_table1818191292017"></a>
8 6 
9| 修饰符 | 函数功能 | 调用方式 |7| 修饰符 | 函数功能 | 调用方式 |
10| --- | --- | --- |8| --- | --- | --- |
@@ -35,9 +33,7 @@ SIMD与SIMT混合编程中涉及的函数类型如下表所示:
35__global__ __vector__ void kernel_name(__gm__ type* param1, __gm__ type* param2, ...);33__global__ __vector__ void kernel_name(__gm__ type* param1, __gm__ type* param2, ...);
36```34```
37 35 
38-关键修饰符说明如下:36+关键修饰符说明如下:<a name="zh-cn_topic_0000002571578013_table4811406341"></a>
39- 
40-<a name="zh-cn_topic_0000002571578013_table4811406341"></a>
41 37 
42| 修饰符 | 作用 | 必需性 |38| 修饰符 | 作用 | 必需性 |
43| --- | --- | --- |39| --- | --- | --- |
@@ -56,9 +52,7 @@ __global__ __vector__ void kernel_name(__gm__ type* param1, __gm__ type* param2,
56kernel_name<<<block_num, dyn_ub_size, stream>>>(args...);52kernel_name<<<block_num, dyn_ub_size, stream>>>(args...);
57```53```
58 54 
59-核函数(Kernel)调用符内的配置参数说明如下:55+核函数(Kernel)调用符内的配置参数说明如下:<a name="zh-cn_topic_0000002571578013_table942016184315"></a>
60- 
61-<a name="zh-cn_topic_0000002571578013_table942016184315"></a>
62 56 
63| 参数 | 类型 | 说明 | 约束 |57| 参数 | 类型 | 说明 | 约束 |
64| --- | --- | --- | --- |58| --- | --- | --- | --- |
@@ -77,9 +71,7 @@ __simt_vf__ __launch_bounds__(MAX_THREAD_COUNT) inline void function_name(
77 type scalar_param, ...);71 type scalar_param, ...);
78```72```
79 73 
80-SIMT VF函数定义中的关键修饰符说明如下:74+SIMT VF函数定义中的关键修饰符说明如下:<a name="zh-cn_topic_0000002571578013_table7661145014492"></a>
81- 
82-<a name="zh-cn_topic_0000002571578013_table7661145014492"></a>
83 75 
84| 修饰符 | 作用 |76| 修饰符 | 作用 |
85| --- | --- |77| --- | --- |
@@ -31,9 +31,8 @@
31 31 
32开发者可以通过[平台信息获取](../../../../../api/Utils-API/platform_info/platform_info.md)接口查询各存储单元的内存空间大小。32开发者可以通过[平台信息获取](../../../../../api/Utils-API/platform_info/platform_info.md)接口查询各存储单元的内存空间大小。
33 33 
34-**各存储单元的最小访问粒度(对齐要求)**34+**各存储单元的最小访问粒度(对齐要求)**<a name="table12348145512210"></a>
35 35 
36-<a name="table12348145512210"></a>
37<table><thead align="left"><tr id="row3348175522218"><th class="cellrowborder" valign="top" width="49.980000000000004%" id="mcps1.1.3.1.1"><p id="p5348145515229"><a name="p5348145515229"></a><a name="p5348145515229"></a>存储单元</p>36<table><thead align="left"><tr id="row3348175522218"><th class="cellrowborder" valign="top" width="49.980000000000004%" id="mcps1.1.3.1.1"><p id="p5348145515229"><a name="p5348145515229"></a><a name="p5348145515229"></a>存储单元</p>
38</th>37</th>
39<th class="cellrowborder" valign="top" width="50.019999999999996%" id="mcps1.1.3.1.2"><p id="p16348125520221"><a name="p16348125520221"></a><a name="p16348125520221"></a>对齐要求</p>38<th class="cellrowborder" valign="top" width="50.019999999999996%" id="mcps1.1.3.1.2"><p id="p16348125520221"><a name="p16348125520221"></a><a name="p16348125520221"></a>对齐要求</p>
@@ -35,9 +35,7 @@
35 35 
36开发者可以通过[平台信息获取](../../../../../api/Utils-API/platform_info/platform_info.md)接口查询各存储单元的内存空间大小。36开发者可以通过[平台信息获取](../../../../../api/Utils-API/platform_info/platform_info.md)接口查询各存储单元的内存空间大小。
37 37 
38-**各存储单元的参数规格**38+**各存储单元的参数规格**<a name="table12348145512210"></a>
39- 
40-<a name="table12348145512210"></a>
41 39 
42| 存储单元名称 | 存储空间大小 | 对齐要求 | 核函数(Kernel)侧常量名称及大小(单位:字节) |40| 存储单元名称 | 存储空间大小 | 对齐要求 | 核函数(Kernel)侧常量名称及大小(单位:字节) |
43| --- | --- | --- | --- |41| --- | --- | --- | --- |
@@ -220,9 +218,7 @@ Fixpipe是NPU将典型操作进行硬化的加速模块,位于AIC内部,配
220 218 
221 同一flagId的计数器最多可以设置15次。219 同一flagId的计数器最多可以设置15次。
222 220 
223- <a id="cross_core_wait_flag_default_values"></a>221+ - **默认值**<a id="cross_core_wait_flag_default_values"></a>
224- 
225- - **默认值**
226 222 
227 - modeId:默认值为0。建议显式配置与CrossCoreSetFlag相同的modeId。223 - modeId:默认值为0。建议显式配置与CrossCoreSetFlag相同的modeId。
228 - pipe:默认值为PIPE\_S。CrossCoreWaitFlag接口提供pipe模板参数,但本架构版本的硬件指令不区分等待的流水,因此pipe取值不影响实际执行效果,CrossCoreWaitFlag会阻塞全部流水的后续指令。224 - pipe:默认值为PIPE\_S。CrossCoreWaitFlag接口提供pipe模板参数,但本架构版本的硬件指令不区分等待的流水,因此pipe取值不影响实际执行效果,CrossCoreWaitFlag会阻塞全部流水的后续指令。
@@ -30,9 +30,8 @@
30 30 
31开发者可以通过[平台信息获取](../../../../../api/Utils-API/platform_info/platform_info.md)接口查询各存储单元的内存空间大小。31开发者可以通过[平台信息获取](../../../../../api/Utils-API/platform_info/platform_info.md)接口查询各存储单元的内存空间大小。
32 32 
33-**各存储单元的最小访问粒度(对齐要求)**33+**各存储单元的最小访问粒度(对齐要求)**<a name="table12348145512210"></a>
34 34 
35-<a name="table12348145512210"></a>
36<table><thead align="left"><tr id="row3348175522218"><th class="cellrowborder" valign="top" width="7.5200000000000005%" id="mcps1.1.4.1.1"><p id="p1680716381143"><a name="p1680716381143"></a><a name="p1680716381143"></a>核</p>35<table><thead align="left"><tr id="row3348175522218"><th class="cellrowborder" valign="top" width="7.5200000000000005%" id="mcps1.1.4.1.1"><p id="p1680716381143"><a name="p1680716381143"></a><a name="p1680716381143"></a>核</p>
37</th>36</th>
38<th class="cellrowborder" valign="top" width="46.22%" id="mcps1.1.4.1.2"><p id="p5348145515229"><a name="p5348145515229"></a><a name="p5348145515229"></a>存储单元</p>37<th class="cellrowborder" valign="top" width="46.22%" id="mcps1.1.4.1.2"><p id="p5348145515229"><a name="p5348145515229"></a><a name="p5348145515229"></a>存储单元</p>
@@ -15,9 +15,8 @@
15- 删除L1 Buffer-\> Global Memory的数据通路。15- 删除L1 Buffer-\> Global Memory的数据通路。
16- SSBuffer,用于AIC和AIV的核间通信。16- SSBuffer,用于AIC和AIV的核间通信。
17- 增加SIMD Register File存储层次,在SIMD程序中,数据从Unified Buffer(UB)搬运到Register进行计算,产生的中间结果可以不用传回UB,直接在寄存器计算。17- 增加SIMD Register File存储层次,在SIMD程序中,数据从Unified Buffer(UB)搬运到Register进行计算,产生的中间结果可以不用传回UB,直接在寄存器计算。
18-- 新增SIMT相关硬件单元。SIMT相关硬件单元介绍如下:18+- 新增SIMT相关硬件单元。SIMT相关硬件单元介绍如下:<a name="table45836437171"></a>
19 19 
20- <a name="table45836437171"></a>
21 <table><thead align="left"><tr id="row1558324391711"><th class="cellrowborder" valign="top" width="18.85%" id="mcps1.1.3.1.1"><p id="p558318431177"><a name="p558318431177"></a><a name="p558318431177"></a>SIMT硬件单元名称</p>20 <table><thead align="left"><tr id="row1558324391711"><th class="cellrowborder" valign="top" width="18.85%" id="mcps1.1.3.1.1"><p id="p558318431177"><a name="p558318431177"></a><a name="p558318431177"></a>SIMT硬件单元名称</p>
22 </th>21 </th>
23 <th class="cellrowborder" valign="top" width="81.15%" id="mcps1.1.3.1.2"><p id="p3583443191718"><a name="p3583443191718"></a><a name="p3583443191718"></a>说明</p>22 <th class="cellrowborder" valign="top" width="81.15%" id="mcps1.1.3.1.2"><p id="p3583443191718"><a name="p3583443191718"></a><a name="p3583443191718"></a>说明</p>
@@ -79,9 +78,7 @@
79 78 
80开发者可以通过[平台信息获取](../../../../../api/Utils-API/platform_info/platform_info.md)接口查询各存储单元的内存空间大小。79开发者可以通过[平台信息获取](../../../../../api/Utils-API/platform_info/platform_info.md)接口查询各存储单元的内存空间大小。
81 80 
82-**各存储单元的参数规格**81+**各存储单元的参数规格**<a name="table12348145512210"></a>
83- 
84-<a name="table12348145512210"></a>
85 82 
86| 存储单元名称 | 存储空间大小 | 对齐要求 | 核函数(Kernel)侧常量名称及大小(单位:字节) |83| 存储单元名称 | 存储空间大小 | 对齐要求 | 核函数(Kernel)侧常量名称及大小(单位:字节) |
87| --- | --- | --- | --- |84| --- | --- | --- | --- |
@@ -104,9 +101,7 @@
104| 非对齐搬入寄存器 | 32B | 4 |101| 非对齐搬入寄存器 | 32B | 4 |
105| 非对齐搬入出寄存器 | 32B | 4 |102| 非对齐搬入出寄存器 | 32B | 4 |
106 103 
107-<a id="compile-options-description"></a>104+> [!NOTE]说明<a id="compile-options-description"></a>
108- 
109-> [!NOTE]说明
110>105>
111> UB总容量为256KB,其中部分空间可能用于系统预留或划分为Data Cache。因此,用户可使用的UB空间大小会随编译选项和编程场景变化。106> UB总容量为256KB,其中部分空间可能用于系统预留或划分为Data Cache。因此,用户可使用的UB空间大小会随编译选项和编程场景变化。
112>107>
@@ -123,9 +123,8 @@ AI处理器中的计算资源要想发挥强劲算力,必要条件是保证输
123**图6** 存储单元<a name="fig2032418113610"></a> 123**图6** 存储单元<a name="fig2032418113610"></a>
124![](../../../figures/storage_unit.png "存储单元")124![](../../../figures/storage_unit.png "存储单元")
125 125 
126-**表1** 存储单元介绍126+**表1** 存储单元介绍<a name="table1692510612218"></a>
127 127 
128-<a name="table1692510612218"></a>
129<table><thead align="left"><tr id="row1692576132110"><th class="cellrowborder" valign="top" width="30.11%" id="mcps1.2.3.1.1"><p id="p189252069218"><a name="p189252069218"></a><a name="p189252069218"></a>存储单元</p>128<table><thead align="left"><tr id="row1692576132110"><th class="cellrowborder" valign="top" width="30.11%" id="mcps1.2.3.1.1"><p id="p189252069218"><a name="p189252069218"></a><a name="p189252069218"></a>存储单元</p>
130</th>129</th>
131<th class="cellrowborder" valign="top" width="69.89%" id="mcps1.2.3.1.2"><p id="p17925161215"><a name="p17925161215"></a><a name="p17925161215"></a>描述</p>130<th class="cellrowborder" valign="top" width="69.89%" id="mcps1.2.3.1.2"><p id="p17925161215"><a name="p17925161215"></a><a name="p17925161215"></a>描述</p>
@@ -165,9 +164,8 @@ AI处理器中的计算资源要想发挥强劲算力,必要条件是保证输
165</tbody>164</tbody>
166</table>165</table>
167 166 
168-**表2** 搬运单元介绍167+**表2** 搬运单元介绍<a name="table288493152012"></a>
169 168 
170-<a name="table288493152012"></a>
171<table><thead align="left"><tr id="row19884103142012"><th class="cellrowborder" valign="top" width="30.049999999999997%" id="mcps1.2.3.1.1"><p id="p288473114204"><a name="p288473114204"></a><a name="p288473114204"></a>搬运单元</p>169<table><thead align="left"><tr id="row19884103142012"><th class="cellrowborder" valign="top" width="30.049999999999997%" id="mcps1.2.3.1.1"><p id="p288473114204"><a name="p288473114204"></a><a name="p288473114204"></a>搬运单元</p>
172</th>170</th>
173<th class="cellrowborder" valign="top" width="69.95%" id="mcps1.2.3.1.2"><p id="p8884531162011"><a name="p8884531162011"></a><a name="p8884531162011"></a>描述</p>171<th class="cellrowborder" valign="top" width="69.95%" id="mcps1.2.3.1.2"><p id="p8884531162011"><a name="p8884531162011"></a><a name="p8884531162011"></a>描述</p>
@@ -18,9 +18,8 @@
18 18 
19下表列出了不同类型的算子对上述实现代码的要求。19下表列出了不同类型的算子对上述实现代码的要求。
20 20 
21-**表1** 不同的类型的算子对入图实现代码的要求21+**表1** 不同的类型的算子对入图实现代码的要求<a name="table772183710452"></a>
22 22 
23-<a name="table772183710452"></a>
24<table><thead align="left"><tr id="row1372110374457"><th class="cellrowborder" valign="top" width="35.8%" id="mcps1.2.3.1.1"><p id="p572263714456"><a name="p572263714456"></a><a name="p572263714456"></a>分类</p>23<table><thead align="left"><tr id="row1372110374457"><th class="cellrowborder" valign="top" width="35.8%" id="mcps1.2.3.1.1"><p id="p572263714456"><a name="p572263714456"></a><a name="p572263714456"></a>分类</p>
25</th>24</th>
26<th class="cellrowborder" valign="top" width="64.2%" id="mcps1.2.3.1.2"><p id="p117222376454"><a name="p117222376454"></a><a name="p117222376454"></a>对入图实现代码的要求</p>25<th class="cellrowborder" valign="top" width="64.2%" id="mcps1.2.3.1.2"><p id="p117222376454"><a name="p117222376454"></a><a name="p117222376454"></a>对入图实现代码的要求</p>
@@ -177,9 +177,8 @@ AddKernelInvocationNeo
177 177 
178本节会介绍CMake文件中一些关键环境变量和Cmake命令参数的说明,通常情况下不需要开发者修改,但是这些参数可以帮助开发者更好的理解编译原理,方便有能力的开发者对Cmake进行定制化处理。178本节会介绍CMake文件中一些关键环境变量和Cmake命令参数的说明,通常情况下不需要开发者修改,但是这些参数可以帮助开发者更好的理解编译原理,方便有能力的开发者对Cmake进行定制化处理。
179 179 
180-**表1** 环境变量说明180+**表1** 环境变量说明<a name="table8269104175818"></a>
181 181 
182-<a name="table8269104175818"></a>
183<table><thead align="left"><tr id="row126924116583"><th class="cellrowborder" valign="top" width="24.54%" id="mcps1.2.3.1.1"><p id="p182692415584"><a name="p182692415584"></a><a name="p182692415584"></a>环境变量</p>182<table><thead align="left"><tr id="row126924116583"><th class="cellrowborder" valign="top" width="24.54%" id="mcps1.2.3.1.1"><p id="p182692415584"><a name="p182692415584"></a><a name="p182692415584"></a>环境变量</p>
184</th>183</th>
185<th class="cellrowborder" valign="top" width="75.46000000000001%" id="mcps1.2.3.1.2"><p id="p12269124113588"><a name="p12269124113588"></a><a name="p12269124113588"></a>配置说明</p>184<th class="cellrowborder" valign="top" width="75.46000000000001%" id="mcps1.2.3.1.2"><p id="p12269124113588"><a name="p12269124113588"></a><a name="p12269124113588"></a>配置说明</p>
@@ -219,9 +218,8 @@ AddKernelInvocationNeo
219</tbody>218</tbody>
220</table>219</table>
221 220 
222-**表2** Cmake命令语法说明221+**表2** Cmake命令语法说明<a name="table481718169817"></a>
223 222 
224-<a name="table481718169817"></a>
225<table><thead align="left"><tr id="row1981751617812"><th class="cellrowborder" valign="top" width="24.67%" id="mcps1.2.3.1.1"><p id="p188171016288"><a name="p188171016288"></a><a name="p188171016288"></a>Cmake命令</p>223<table><thead align="left"><tr id="row1981751617812"><th class="cellrowborder" valign="top" width="24.67%" id="mcps1.2.3.1.1"><p id="p188171016288"><a name="p188171016288"></a><a name="p188171016288"></a>Cmake命令</p>
226</th>224</th>
227<th class="cellrowborder" valign="top" width="75.33%" id="mcps1.2.3.1.2"><p id="p481751615812"><a name="p481751615812"></a><a name="p481751615812"></a>语法说明</p>225<th class="cellrowborder" valign="top" width="75.33%" id="mcps1.2.3.1.2"><p id="p481751615812"><a name="p481751615812"></a><a name="p481751615812"></a>语法说明</p>
@@ -386,9 +384,8 @@ bash run.sh --run-mode=npu --soc-version=<soc_version> --install-path=<install_
386bash run.sh -r npu -v <soc_version> -i <install_path> -b Debug -p <install-prefix>384bash run.sh -r npu -v <soc_version> -i <install_path> -b Debug -p <install-prefix>
387```385```
388 386 
389-**表3** 脚本参数介绍387+**表3** 脚本参数介绍<a name="table98393011180"></a>
390 388 
391-<a name="table98393011180"></a>
392<table><thead align="left"><tr id="row98396051814"><th class="cellrowborder" valign="top" width="16.14%" id="mcps1.2.4.1.1"><p id="p283916071814"><a name="p283916071814"></a><a name="p283916071814"></a>参数名</p>389<table><thead align="left"><tr id="row98396051814"><th class="cellrowborder" valign="top" width="16.14%" id="mcps1.2.4.1.1"><p id="p283916071814"><a name="p283916071814"></a><a name="p283916071814"></a>参数名</p>
393</th>390</th>
394<th class="cellrowborder" valign="top" width="11.690000000000001%" id="mcps1.2.4.1.2"><p id="p1872572854511"><a name="p1872572854511"></a><a name="p1872572854511"></a>参数简写</p>391<th class="cellrowborder" valign="top" width="11.690000000000001%" id="mcps1.2.4.1.2"><p id="p1872572854511"><a name="p1872572854511"></a><a name="p1872572854511"></a>参数简写</p>
@@ -95,9 +95,7 @@
95 msobjdump --list-elf <elf_file>95 msobjdump --list-elf <elf_file>
96 ```96 ```
97 97 
98- **表3** 参数说明98+ **表3** 参数说明<a name="table121952819427"></a>
99- 
100- <a name="table121952819427"></a>
101 99 
102 | 参数(区分大小写) | 可选/必选 | 说明 |100 | 参数(区分大小写) | 可选/必选 | 说明 |
103 | ---- | ---- | ---- |101 | ---- | ---- | ---- |
@@ -133,9 +131,7 @@
133| `FUNCTION_ENTRY` | 算子TilingKey的值。<br><!-- npu="950" id55 -->Ascend 950PR&950DT系列产品:支持<br><!-- end id55 --><!-- npu="A3" id56 -->Atlas A3系列产品:不支持<br><!-- end id56 --><!-- npu="910b" id57 -->Atlas A2系列产品:不支持<br><!-- end id57 --><!-- npu="310b" id58 -->Atlas 200I/500 A2推理产品:不支持<br><!-- end id58 --><!-- npu="310p" id59 -->Atlas 推理系列产品:不支持<br><!-- end id59 --><!-- npu="910" id60 -->Atlas训练系列产品:不支持<!-- end id60 --> | 否 | 不设置`--verbose`,默认打印。 |131| `FUNCTION_ENTRY` | 算子TilingKey的值。<br><!-- npu="950" id55 -->Ascend 950PR&950DT系列产品:支持<br><!-- end id55 --><!-- npu="A3" id56 -->Atlas A3系列产品:不支持<br><!-- end id56 --><!-- npu="910b" id57 -->Atlas A2系列产品:不支持<br><!-- end id57 --><!-- npu="310b" id58 -->Atlas 200I/500 A2推理产品:不支持<br><!-- end id58 --><!-- npu="310p" id59 -->Atlas 推理系列产品:不支持<br><!-- end id59 --><!-- npu="910" id60 -->Atlas训练系列产品:不支持<!-- end id60 --> | 否 | 不设置`--verbose`,默认打印。 |
134| `elf header infos` | 包括ELF Header、Section Headers、Key to Flags、Program Headers、Symbol表等信息。 | 否 | 设置`--verbose`,开启全量打印。 |132| `elf header infos` | 包括ELF Header、Section Headers、Key to Flags、Program Headers、Symbol表等信息。 | 否 | 设置`--verbose`,开启全量打印。 |
135 133 
136-**表6** 核函数(Kernel)类型信息134+**表6** 核函数(Kernel)类型信息<a name="table187419221164"></a>
137- 
138-<a name="table187419221164"></a>
139 135 
140| KERNEL_TYPE | 说明 |136| KERNEL_TYPE | 说明 |
141| ---- | ---- |137| ---- | ---- |
@@ -133,9 +133,7 @@ optype_collector工具用于采集内置/自定义算子包中指定AI处理器
133 | ---- | ---- | ---- |133 | ---- | ---- | ---- |
134 | `--detect-conflicts {soc_version}` | 必选 | 检测指定AI处理器版本下自定义算子OpType是否与CANN内置算子OpType或其他自定义算子OpType重名。`{soc_version}`表示待检测的AI处理器型号,可填写CANN支持的产品AI处理器型号。 |134 | `--detect-conflicts {soc_version}` | 必选 | 检测指定AI处理器版本下自定义算子OpType是否与CANN内置算子OpType或其他自定义算子OpType重名。`{soc_version}`表示待检测的AI处理器型号,可填写CANN支持的产品AI处理器型号。 |
135 135 
136-## 返回值<a name="section_optype_collector_return_code"></a>136+## 返回值<a name="section_optype_collector_return_code"></a><a name="table_optype_collector_return_code"></a>
137- 
138-<a name="table_optype_collector_return_code"></a>
139 137 
140| 返回值 | 说明 |138| 返回值 | 说明 |
141| ---- | ---- |139| ---- | ---- |
@@ -145,9 +143,7 @@ optype_collector工具用于采集内置/自定义算子包中指定AI处理器
145 143 
146## 输出说明<a name="section_optype_collector_output"></a>144## 输出说明<a name="section_optype_collector_output"></a>
147 145 
148-工具执行后会根据命令模式输出如下信息。146+工具执行后会根据命令模式输出如下信息。<a name="table_optype_collector_output"></a>
149- 
150-<a name="table_optype_collector_output"></a>
151 147 
152| 输出项 | 说明 |148| 输出项 | 说明 |
153| ---- | ---- |149| ---- | ---- |
@@ -163,9 +163,7 @@ AI Core SIMT的基本编译流程如下:Host代码使用Host编译器编译成
163 163 
164更多的编译命令和用法可以参见[《毕昇编译器用户指南》](https://www.hiascend.com/document/redirect/CannCommunityBiSheng)。164更多的编译命令和用法可以参见[《毕昇编译器用户指南》](https://www.hiascend.com/document/redirect/CannCommunityBiSheng)。
165 165 
166-<a id="note-reference"></a>166+>[!NOTE]说明<a id="note-reference"></a>
167- 
168->[!NOTE]说明
169><!-- npu="910b,910,310p,310b" id1 -->167><!-- npu="910b,910,310p,310b" id1 -->
170>- 针对如下产品:在安装AI处理器的服务器执行`npu-smi info`命令进行查询,获取Name信息。实际配置值为AscendName,例如Name取值为xxxyy,实际配置值为Ascendxxxyy。<br><br>168>- 针对如下产品:在安装AI处理器的服务器执行`npu-smi info`命令进行查询,获取Name信息。实际配置值为AscendName,例如Name取值为xxxyy,实际配置值为Ascendxxxyy。<br><br>
171> <!-- npu="910b" id2 -->169> <!-- npu="910b" id2 -->
@@ -84,9 +84,8 @@ bisheng hello_world.aicpu.o main.asc.o -o demo
84 84 
85## AI CPU算子常用编译选项<a name="section345885113142"></a>85## AI CPU算子常用编译选项<a name="section345885113142"></a>
86 86 
87-AI CPU算子常用的编译选项说明如下:87+AI CPU算子常用的编译选项说明如下:<a name="table9126181131320"></a>
88 88 
89-<a name="table9126181131320"></a>
90<table><thead align="left"><tr id="row312711101316"><th class="cellrowborder" valign="top" width="33.63636363636363%" id="mcps1.1.4.1.1"><p id="p71271711201318"><a name="p71271711201318"></a><a name="p71271711201318"></a><strong id="b01279110139"><a name="b01279110139"></a><a name="b01279110139"></a>选项</strong></p>89<table><thead align="left"><tr id="row312711101316"><th class="cellrowborder" valign="top" width="33.63636363636363%" id="mcps1.1.4.1.1"><p id="p71271711201318"><a name="p71271711201318"></a><a name="p71271711201318"></a><strong id="b01279110139"><a name="b01279110139"></a><a name="b01279110139"></a>选项</strong></p>
91</th>90</th>
92<th class="cellrowborder" valign="top" width="9.676767676767676%" id="mcps1.1.4.1.2"><p id="p1212711115131"><a name="p1212711115131"></a><a name="p1212711115131"></a><strong id="b101271011101310"><a name="b101271011101310"></a><a name="b101271011101310"></a>是否必需</strong></p>91<th class="cellrowborder" valign="top" width="9.676767676767676%" id="mcps1.1.4.1.2"><p id="p1212711115131"><a name="p1212711115131"></a><a name="p1212711115131"></a><strong id="b101271011101310"><a name="b101271011101310"></a><a name="b101271011101310"></a>是否必需</strong></p>
@@ -6,9 +6,8 @@ Ascend C算子调试的整体方案如下:开发者通过调用Ascend C类库
6 6 
7具体的调试调优方法和使用的工具列表如下:7具体的调试调优方法和使用的工具列表如下:
8 8 
9-**表1** 调试调优方法和使用的工具列表9+**表1** 调试调优方法和使用的工具列表<a name="table3879116815"></a>
10 10 
11-<a name="table3879116815"></a>
12<table><thead align="left"><tr id="row128713114816"><th class="cellrowborder" valign="top" width="10.41104110411041%" id="mcps1.2.4.1.1"><p id="p98701186"><a name="p98701186"></a><a name="p98701186"></a>分类</p>11<table><thead align="left"><tr id="row128713114816"><th class="cellrowborder" valign="top" width="10.41104110411041%" id="mcps1.2.4.1.1"><p id="p98701186"><a name="p98701186"></a><a name="p98701186"></a>分类</p>
13</th>12</th>
14<th class="cellrowborder" valign="top" width="15.541554155415543%" id="mcps1.2.4.1.2"><p id="p16871313811"><a name="p16871313811"></a><a name="p16871313811"></a>子分类</p>13<th class="cellrowborder" valign="top" width="15.541554155415543%" id="mcps1.2.4.1.2"><p id="p16871313811"><a name="p16871313811"></a><a name="p16871313811"></a>子分类</p>
@@ -48,8 +48,7 @@
48 └──visualize_data.bin # MindStudio Insight呈现文件48 └──visualize_data.bin # MindStudio Insight呈现文件
49 ```49 ```
50 50
51-**表2 msopprof文件介绍**51+**表2 msopprof文件介绍**<a name="table244174535419"></a>
52-<a name="table244174535419"></a>
53<table><thead align="left"><tr id="row17451145145411"><th class="cellrowborder" valign="top" width="24.54%" id="mcps1.2.3.1.1"><p id="p18451345115414"><a name="p18451345115414"></a><a name="p18451345115414"></a>名称</p>52<table><thead align="left"><tr id="row17451145145411"><th class="cellrowborder" valign="top" width="24.54%" id="mcps1.2.3.1.1"><p id="p18451345115414"><a name="p18451345115414"></a><a name="p18451345115414"></a>名称</p>
54</th>53</th>
55<th class="cellrowborder" valign="top" width="75.46000000000001%" id="mcps1.2.3.1.2"><p id="p4451245115415"><a name="p4451245115415"></a><a name="p4451245115415"></a>说明</p>54<th class="cellrowborder" valign="top" width="75.46000000000001%" id="mcps1.2.3.1.2"><p id="p4451245115415"><a name="p4451245115415"></a><a name="p4451245115415"></a>说明</p>
@@ -160,9 +159,8 @@
160 159 
161 其中仿真编译所依赖的库介绍如下,开启仿真编译时,需要优先链接,确保优先使用仿真库的符号,防止出现运行时coredump等异常情况。160 其中仿真编译所依赖的库介绍如下,开启仿真编译时,需要优先链接,确保优先使用仿真库的符号,防止出现运行时coredump等异常情况。
162 161 
163- **表3 仿真编译依赖库介绍**162+ **表3 仿真编译依赖库介绍**<a name="table84131311195116"></a>
164 163 
165- <a name="table84131311195116"></a>
166 <table><thead align="left"><tr id="row4413911115113"><th class="cellrowborder" valign="top" width="45.97%" id="mcps1.2.3.1.1"><p id="p1841381120511"><a name="p1841381120511"></a><a name="p1841381120511"></a>名称</p>164 <table><thead align="left"><tr id="row4413911115113"><th class="cellrowborder" valign="top" width="45.97%" id="mcps1.2.3.1.1"><p id="p1841381120511"><a name="p1841381120511"></a><a name="p1841381120511"></a>名称</p>
167 </th>165 </th>
168 <th class="cellrowborder" valign="top" width="54.03%" id="mcps1.2.3.1.2"><p id="p15413191116515"><a name="p15413191116515"></a><a name="p15413191116515"></a>作用描述</p>166 <th class="cellrowborder" valign="top" width="54.03%" id="mcps1.2.3.1.2"><p id="p15413191116515"><a name="p15413191116515"></a><a name="p15413191116515"></a>作用描述</p>
@@ -31,9 +31,7 @@ SIMT会提供内置的宏方便用户编写程序。预定义宏一节着重介
31 31 
32函数执行空间限定符(Function Execution Space Qualifier)指示函数是在Host侧执行还是在Device侧执行,以及能被调用的空间范围。32函数执行空间限定符(Function Execution Space Qualifier)指示函数是在Host侧执行还是在Device侧执行,以及能被调用的空间范围。
33 33 
34-**表1** 函数执行空间限定符概览34+**表1** 函数执行空间限定符概览<a name="table121121062614"></a>
35- 
36-<a name="table121121062614"></a>
37 35 
38| 函数执行空间限定符 | 执行空间(Host) | 执行空间(Device) | 允许调用函数空间(Host) | 允许调用函数空间(Device) |36| 函数执行空间限定符 | 执行空间(Host) | 执行空间(Device) | 允许调用函数空间(Host) | 允许调用函数空间(Device) |
39| --- | --- | --- | --- | --- |37| --- | --- | --- | --- | --- |
@@ -70,9 +68,7 @@ SIMT会提供内置的宏方便用户编写程序。预定义宏一节着重介
70 68 
71 动态内存的实际内存大小需要在核函数(Kernel)启动时配置,具体内容请参考[核函数(Kernel)配置](#section97005415463)。69 动态内存的实际内存大小需要在核函数(Kernel)启动时配置,具体内容请参考[核函数(Kernel)配置](#section97005415463)。
72 70 
73-## 内置常量<a name="section784531219338"></a>71+## 内置常量<a name="section784531219338"></a><a name="table878553753515"></a>
74- 
75-<a name="table878553753515"></a>
76 72 
77| 常量名 | 取值 | 功能 |73| 常量名 | 取值 | 功能 |
78| --- | --- | --- |74| --- | --- | --- |
@@ -153,9 +149,7 @@ SIMT会提供内置的宏方便用户编写程序。预定义宏一节着重介
153 149 
154目前提供了一系列适用于Device侧的数据类型,包括标量和短向量。短向量是由多个元素组成的简单向量。以下数据类型是编译器默认支持的数据类型,使用时无需引入头文件。150目前提供了一系列适用于Device侧的数据类型,包括标量和短向量。短向量是由多个元素组成的简单向量。以下数据类型是编译器默认支持的数据类型,使用时无需引入头文件。
155 151 
156-**表2** 标量数据类型152+**表2** 标量数据类型<a name="table114921244133211"></a>
157- 
158-<a name="table114921244133211"></a>
159 153 
160| 类型 | 数据类型 | 描述 | Size(bit) | 取值范围 |154| 类型 | 数据类型 | 描述 | Size(bit) | 取值范围 |
161| --- | --- | --- | --- | --- |155| --- | --- | --- | --- | --- |
@@ -175,9 +169,7 @@ SIMT会提供内置的宏方便用户编写程序。预定义宏一节着重介
175| 浮点型 | bfloat16_t | 符号位宽1,指数位宽8,尾数位宽7 | 16 | [2<sup>120</sup> - 2<sup>128</sup>, 2<sup>128</sup> - 2<sup>120</sup>] |169| 浮点型 | bfloat16_t | 符号位宽1,指数位宽8,尾数位宽7 | 16 | [2<sup>120</sup> - 2<sup>128</sup>, 2<sup>128</sup> - 2<sup>120</sup>] |
176| 浮点型 | float | 符号位宽1,指数位宽8,尾数位宽23 | 32 | [2<sup>104</sup> - 2<sup>128</sup>, 2<sup>128</sup> - 2<sup>104</sup>] |170| 浮点型 | float | 符号位宽1,指数位宽8,尾数位宽23 | 32 | [2<sup>104</sup> - 2<sup>128</sup>, 2<sup>128</sup> - 2<sup>104</sup>] |
177 171 
178-短向量数据类型分为Vector X2、Vector X3、Vector X4,表示一个短向量变量有2、3、4个元素,当前支持的类型分布如下:172+短向量数据类型分为Vector X2、Vector X3、Vector X4,表示一个短向量变量有2、3、4个元素,当前支持的类型分布如下:<a name="table49986503175"></a>
179- 
180-<a name="table49986503175"></a>
181 173 
182| 元素数据类型 | Vector X2 | Vector X3 | Vector X4 |174| 元素数据类型 | Vector X2 | Vector X3 | Vector X4 |
183| --- | --- | --- | --- |175| --- | --- | --- | --- |
@@ -200,9 +192,7 @@ SIMT会提供内置的宏方便用户编写程序。预定义宏一节着重介
200| 浮点型,1符号位,8指数位,7尾数位 | bfloat16x2_t | - | - |192| 浮点型,1符号位,8指数位,7尾数位 | bfloat16x2_t | - | - |
201| 浮点型,1符号位,8指数位,23尾数位 | float2 | float3 | float4 |193| 浮点型,1符号位,8指数位,23尾数位 | float2 | float3 | float4 |
202 194 
203-**表3** 短向量数据类型195+**表3** 短向量数据类型<a name="table13856144241"></a>
204- 
205-<a name="table13856144241"></a>
206 196 
207| 数据类型 | 内存大小(字节) | 地址对齐(字节) |197| 数据类型 | 内存大小(字节) | 地址对齐(字节) |
208| --- | --- | --- |198| --- | --- | --- |
@@ -226,9 +216,7 @@ SIMT会提供内置的宏方便用户编写程序。预定义宏一节着重介
226 216 
227SIMT编程提供了一系列运算符,用于执行数学运算。以下是支持的运算符列表。217SIMT编程提供了一系列运算符,用于执行数学运算。以下是支持的运算符列表。
228 218 
229-**表4** SIMT编程支持的运算符列表219+**表4** SIMT编程支持的运算符列表<a name="table11745172843710"></a>
230- 
231-<a name="table11745172843710"></a>
232 220 
233| 类别 | 运算符 | bool | int8_t/uint8_t/int16_t/uint16_t/int32_t/uint32_t/int64_t/uint64_t | half/bfloat16_t/float | half2/bfloat16x2_t | hifloat8_t |221| 类别 | 运算符 | bool | int8_t/uint8_t/int16_t/uint16_t/int32_t/uint32_t/int64_t/uint64_t | half/bfloat16_t/float | half2/bfloat16x2_t | hifloat8_t |
234| --- | --- | --- | --- | --- | --- | --- |222| --- | --- | --- | --- | --- | --- | --- |
@@ -328,10 +316,8 @@ add_custom<<<blocks_per_grid, threads_per_block, dyn_ubuf_size, stream>>>(x, y,
328 316
329 函数标记宏,在核函数(Kernel)上可选配置,用于指定核函数(Kernel)启动的最大线程数。最大线程数决定了每个线程可分配的寄存器数量,具体对应关系请见下表,寄存器用于存储线程中的局部变量,若局部变量的个数超出寄存器个数,容易出现寄存器溢出等问题。建议最大线程数与启动核函数(Kernel)时的dim3线程数保持一致。317 函数标记宏,在核函数(Kernel)上可选配置,用于指定核函数(Kernel)启动的最大线程数。最大线程数决定了每个线程可分配的寄存器数量,具体对应关系请见下表,寄存器用于存储线程中的局部变量,若局部变量的个数超出寄存器个数,容易出现寄存器溢出等问题。建议最大线程数与启动核函数(Kernel)时的dim3线程数保持一致。
330 318 
331- **表5** \_\_launch\_bounds\_\_(N)与每个线程可用的寄存器个数的关系319+ **表5** \_\_launch\_bounds\_\_(N)与每个线程可用的寄存器个数的关系<a name="table1715318510594"></a>
332 320
333- <a name="table1715318510594"></a>
334- 
335 | N | 每个线程可用寄存器个数 |321 | N | 每个线程可用寄存器个数 |
336 | --- | --- |322 | --- | --- |
337 | 1025~2048 | 16 |323 | 1025~2048 | 16 |
@@ -2,9 +2,8 @@
2 2 
3Mask用于控制矢量计算中参与计算的元素个数,支持以下工作模式及配置方式:3Mask用于控制矢量计算中参与计算的元素个数,支持以下工作模式及配置方式:
4 4 
5-**表1** Mask工作模式5+**表1** Mask工作模式<a name="zh-cn_topic_0000002267504584_table414483923116"></a>
6 6 
7-<a name="zh-cn_topic_0000002267504584_table414483923116"></a>
8<table><thead align="left"><tr id="zh-cn_topic_0000002267504584_row51452039163114"><th class="cellrowborder" valign="top" width="14.14%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002267504584_p814513911319"><a name="zh-cn_topic_0000002267504584_p814513911319"></a><a name="zh-cn_topic_0000002267504584_p814513911319"></a>工作模式</p>7<table><thead align="left"><tr id="zh-cn_topic_0000002267504584_row51452039163114"><th class="cellrowborder" valign="top" width="14.14%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002267504584_p814513911319"><a name="zh-cn_topic_0000002267504584_p814513911319"></a><a name="zh-cn_topic_0000002267504584_p814513911319"></a>工作模式</p>
9</th>8</th>
10<th class="cellrowborder" valign="top" width="85.86%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002267504584_p14145239103115"><a name="zh-cn_topic_0000002267504584_p14145239103115"></a><a name="zh-cn_topic_0000002267504584_p14145239103115"></a>说明</p>9<th class="cellrowborder" valign="top" width="85.86%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002267504584_p14145239103115"><a name="zh-cn_topic_0000002267504584_p14145239103115"></a><a name="zh-cn_topic_0000002267504584_p14145239103115"></a>说明</p>
@@ -28,9 +27,8 @@ Mask用于控制矢量计算中参与计算的元素个数,支持以下工作
28</tbody>27</tbody>
29</table>28</table>
30 29 
31-**表2** Mask配置方式30+**表2** Mask配置方式<a name="zh-cn_topic_0000002267504584_table642464733119"></a>
32 31 
33-<a name="zh-cn_topic_0000002267504584_table642464733119"></a>
34<table><thead align="left"><tr id="zh-cn_topic_0000002267504584_row8425134710316"><th class="cellrowborder" valign="top" width="14.12%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002267504584_p194258470311"><a name="zh-cn_topic_0000002267504584_p194258470311"></a><a name="zh-cn_topic_0000002267504584_p194258470311"></a>配置方式</p>32<table><thead align="left"><tr id="zh-cn_topic_0000002267504584_row8425134710316"><th class="cellrowborder" valign="top" width="14.12%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000002267504584_p194258470311"><a name="zh-cn_topic_0000002267504584_p194258470311"></a><a name="zh-cn_topic_0000002267504584_p194258470311"></a>配置方式</p>
35</th>33</th>
36<th class="cellrowborder" valign="top" width="85.88%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002267504584_p442513473313"><a name="zh-cn_topic_0000002267504584_p442513473313"></a><a name="zh-cn_topic_0000002267504584_p442513473313"></a>说明</p>34<th class="cellrowborder" valign="top" width="85.88%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000002267504584_p442513473313"><a name="zh-cn_topic_0000002267504584_p442513473313"></a><a name="zh-cn_topic_0000002267504584_p442513473313"></a>说明</p>
@@ -52,9 +50,8 @@ Mask用于控制矢量计算中参与计算的元素个数,支持以下工作
52 50 
53Mask操作的使用方式如下:51Mask操作的使用方式如下:
54 52 
55-**表3** Mask操作的使用方式53+**表3** Mask操作的使用方式<a name="zh-cn_topic_0000002267504584_table1957843418427"></a>
56 54 
57-<a name="zh-cn_topic_0000002267504584_table1957843418427"></a>
58<table><thead align="left"><tr id="zh-cn_topic_0000002267504584_row657817341420"><th class="cellrowborder" valign="top" width="8.58%" id="mcps1.2.5.1.1"><p id="zh-cn_topic_0000002267504584_p19578634194213"><a name="zh-cn_topic_0000002267504584_p19578634194213"></a><a name="zh-cn_topic_0000002267504584_p19578634194213"></a>配置方式</p>55<table><thead align="left"><tr id="zh-cn_topic_0000002267504584_row657817341420"><th class="cellrowborder" valign="top" width="8.58%" id="mcps1.2.5.1.1"><p id="zh-cn_topic_0000002267504584_p19578634194213"><a name="zh-cn_topic_0000002267504584_p19578634194213"></a><a name="zh-cn_topic_0000002267504584_p19578634194213"></a>配置方式</p>
59</th>56</th>
60<th class="cellrowborder" valign="top" width="8.37%" id="mcps1.2.5.1.2"><p id="zh-cn_topic_0000002267504584_p165791734184211"><a name="zh-cn_topic_0000002267504584_p165791734184211"></a><a name="zh-cn_topic_0000002267504584_p165791734184211"></a>工作模式</p>57<th class="cellrowborder" valign="top" width="8.37%" id="mcps1.2.5.1.2"><p id="zh-cn_topic_0000002267504584_p165791734184211"><a name="zh-cn_topic_0000002267504584_p165791734184211"></a><a name="zh-cn_topic_0000002267504584_p165791734184211"></a>工作模式</p>
@@ -18,9 +18,7 @@ AI Core的抽象硬件架构可以分为[**计算单元**](#section_compute_unit
18 18 
19AI Core中的计算单元主要包括Scalar、Vector和Cube三类。19AI Core中的计算单元主要包括Scalar、Vector和Cube三类。
20 20 
21-**表1** 计算单元21+**表1** 计算单元<a name="zh-cn_topic_0000001588832845_table14884154362"></a>
22- 
23-<a name="zh-cn_topic_0000001588832845_table14884154362"></a>
24 22 
25| 组件名称 | 组件功能 | 对应编程概念 |23| 组件名称 | 组件功能 | 对应编程概念 |
26| --- | --- | --- |24| --- | --- | --- |
@@ -405,9 +405,8 @@ __global__ __mix__ void mix_kernel(__gm__ float* x, __gm__ float* y, __gm__ floa
405## 支持的API范围<a name="section2633193623711"></a>405## 支持的API范围<a name="section2633193623711"></a>
406 406 
407<!-- npu="310p" id1 -->407<!-- npu="310p" id1 -->
408-**表1** 针对Atlas推理系列产品AI Core,支持的API范围408+**表1** 针对Atlas推理系列产品AI Core,支持的API范围<a name="table56285428438"></a>
409 409 
410-<a name="table56285428438"></a>
411<table><thead align="left"><tr id="row46697425436"><th class="cellrowborder" valign="top" width="30.28%" id="mcps1.2.3.1.1"><p id="p866911429430"><a name="p866911429430"></a><a name="p866911429430"></a>接口分类</p>410<table><thead align="left"><tr id="row46697425436"><th class="cellrowborder" valign="top" width="30.28%" id="mcps1.2.3.1.1"><p id="p866911429430"><a name="p866911429430"></a><a name="p866911429430"></a>接口分类</p>
412</th>411</th>
413<th class="cellrowborder" valign="top" width="69.72%" id="mcps1.2.3.1.2"><p id="p186691442174319"><a name="p186691442174319"></a><a name="p186691442174319"></a>接口名称</p>412<th class="cellrowborder" valign="top" width="69.72%" id="mcps1.2.3.1.2"><p id="p186691442174319"><a name="p186691442174319"></a><a name="p186691442174319"></a>接口名称</p>
@@ -515,9 +514,8 @@ __global__ __mix__ void mix_kernel(__gm__ float* x, __gm__ float* y, __gm__ floa
515<!-- end id1 -->514<!-- end id1 -->
516 515 
517<!-- npu="910b" id2 -->516<!-- npu="910b" id2 -->
518-**表2** 针对Atlas A2系列产品,支持的API范围517+**表2** 针对Atlas A2系列产品,支持的API范围<a name="table1842818536431"></a>
519 518 
520-<a name="table1842818536431"></a>
521<table><thead align="left"><tr id="row104742531436"><th class="cellrowborder" valign="top" width="29.9%" id="mcps1.2.4.1.1"><p id="p10474205364319"><a name="p10474205364319"></a><a name="p10474205364319"></a>接口分类</p>519<table><thead align="left"><tr id="row104742531436"><th class="cellrowborder" valign="top" width="29.9%" id="mcps1.2.4.1.1"><p id="p10474205364319"><a name="p10474205364319"></a><a name="p10474205364319"></a>接口分类</p>
522</th>520</th>
523<th class="cellrowborder" valign="top" width="44.25%" id="mcps1.2.4.1.2"><p id="p1047418533430"><a name="p1047418533430"></a><a name="p1047418533430"></a>接口名称</p>521<th class="cellrowborder" valign="top" width="44.25%" id="mcps1.2.4.1.2"><p id="p1047418533430"><a name="p1047418533430"></a><a name="p1047418533430"></a>接口名称</p>
@@ -718,9 +716,8 @@ __global__ __mix__ void mix_kernel(__gm__ float* x, __gm__ float* y, __gm__ floa
718<!-- end id2 -->716<!-- end id2 -->
719 717 
720<!-- npu="A3" id3 -->718<!-- npu="A3" id3 -->
721-**表3** 针对Atlas A3系列产品,支持的API范围719+**表3** 针对Atlas A3系列产品,支持的API范围<a name="table167051848128"></a>
722 720 
723-<a name="table167051848128"></a>
724<table><thead align="left"><tr id="row07054481922"><th class="cellrowborder" valign="top" width="29.630000000000003%" id="mcps1.2.4.1.1"><p id="p135155581014"><a name="p135155581014"></a><a name="p135155581014"></a>接口分类</p>721<table><thead align="left"><tr id="row07054481922"><th class="cellrowborder" valign="top" width="29.630000000000003%" id="mcps1.2.4.1.1"><p id="p135155581014"><a name="p135155581014"></a><a name="p135155581014"></a>接口分类</p>
725</th>722</th>
726<th class="cellrowborder" valign="top" width="44.519999999999996%" id="mcps1.2.4.1.2"><p id="p12351955181013"><a name="p12351955181013"></a><a name="p12351955181013"></a>接口名称</p>723<th class="cellrowborder" valign="top" width="44.519999999999996%" id="mcps1.2.4.1.2"><p id="p12351955181013"><a name="p12351955181013"></a><a name="p12351955181013"></a>接口名称</p>
@@ -934,9 +931,8 @@ __global__ __mix__ void mix_kernel(__gm__ float* x, __gm__ float* y, __gm__ floa
934<!-- end id3 -->931<!-- end id3 -->
935 932 
936<!-- npu="950" id4 -->933<!-- npu="950" id4 -->
937-**表4** 针对Ascend 950PR&950DT系列产品,支持的API范围934+**表4** 针对Ascend 950PR&950DT系列产品,支持的API范围<a name="table1798673512413"></a>
938 935 
939-<a name="table1798673512413"></a>
940<table><thead align="left"><tr id="row109864353413"><th class="cellrowborder" valign="top" width="29.630000000000003%" id="mcps1.2.4.1.1"><p id="p1098613524117"><a name="p1098613524117"></a><a name="p1098613524117"></a>接口分类</p>936<table><thead align="left"><tr id="row109864353413"><th class="cellrowborder" valign="top" width="29.630000000000003%" id="mcps1.2.4.1.1"><p id="p1098613524117"><a name="p1098613524117"></a><a name="p1098613524117"></a>接口分类</p>
941</th>937</th>
942<th class="cellrowborder" valign="top" width="44.519999999999996%" id="mcps1.2.4.1.2"><p id="p29871535174117"><a name="p29871535174117"></a><a name="p29871535174117"></a>接口名称</p>938<th class="cellrowborder" valign="top" width="44.519999999999996%" id="mcps1.2.4.1.2"><p id="p29871535174117"><a name="p29871535174117"></a><a name="p29871535174117"></a>接口名称</p>
@@ -185,9 +185,8 @@
185## 支持的API范围<a name="section2633193623711"></a>185## 支持的API范围<a name="section2633193623711"></a>
186 186 
187<!-- npu="310p" id1 -->187<!-- npu="310p" id1 -->
188-**表1** 针对Atlas推理系列产品AI Core,支持的API范围188+**表1** 针对Atlas推理系列产品AI Core,支持的API范围<a name="table56285428438"></a>
189 189 
190-<a name="table56285428438"></a>
191<table><thead align="left"><tr id="row46697425436"><th class="cellrowborder" valign="top" width="30.28%" id="mcps1.2.3.1.1"><p id="p866911429430"><a name="p866911429430"></a><a name="p866911429430"></a>接口分类</p>190<table><thead align="left"><tr id="row46697425436"><th class="cellrowborder" valign="top" width="30.28%" id="mcps1.2.3.1.1"><p id="p866911429430"><a name="p866911429430"></a><a name="p866911429430"></a>接口分类</p>
192</th>191</th>
193<th class="cellrowborder" valign="top" width="69.72%" id="mcps1.2.3.1.2"><p id="p186691442174319"><a name="p186691442174319"></a><a name="p186691442174319"></a>接口名称</p>192<th class="cellrowborder" valign="top" width="69.72%" id="mcps1.2.3.1.2"><p id="p186691442174319"><a name="p186691442174319"></a><a name="p186691442174319"></a>接口名称</p>
@@ -295,9 +294,8 @@
295<!-- end id1 -->294<!-- end id1 -->
296 295 
297<!-- npu="910b" id2 -->296<!-- npu="910b" id2 -->
298-**表2** 针对Atlas A2系列产品,支持的API范围297+**表2** 针对Atlas A2系列产品,支持的API范围<a name="table1842818536431"></a>
299 298 
300-<a name="table1842818536431"></a>
301<table><thead align="left"><tr id="row104742531436"><th class="cellrowborder" valign="top" width="29.9%" id="mcps1.2.4.1.1"><p id="p10474205364319"><a name="p10474205364319"></a><a name="p10474205364319"></a>接口分类</p>299<table><thead align="left"><tr id="row104742531436"><th class="cellrowborder" valign="top" width="29.9%" id="mcps1.2.4.1.1"><p id="p10474205364319"><a name="p10474205364319"></a><a name="p10474205364319"></a>接口分类</p>
302</th>300</th>
303<th class="cellrowborder" valign="top" width="44.25%" id="mcps1.2.4.1.2"><p id="p1047418533430"><a name="p1047418533430"></a><a name="p1047418533430"></a>接口名称</p>301<th class="cellrowborder" valign="top" width="44.25%" id="mcps1.2.4.1.2"><p id="p1047418533430"><a name="p1047418533430"></a><a name="p1047418533430"></a>接口名称</p>
@@ -498,9 +496,8 @@
498<!-- end id2 -->496<!-- end id2 -->
499 497 
500<!-- npu="A3" id3 -->498<!-- npu="A3" id3 -->
501-**表3** 针对Atlas A3系列产品,支持的API范围499+**表3** 针对Atlas A3系列产品,支持的API范围<a name="table167051848128"></a>
502 500 
503-<a name="table167051848128"></a>
504<table><thead align="left"><tr id="row07054481922"><th class="cellrowborder" valign="top" width="29.630000000000003%" id="mcps1.2.4.1.1"><p id="p135155581014"><a name="p135155581014"></a><a name="p135155581014"></a>接口分类</p>501<table><thead align="left"><tr id="row07054481922"><th class="cellrowborder" valign="top" width="29.630000000000003%" id="mcps1.2.4.1.1"><p id="p135155581014"><a name="p135155581014"></a><a name="p135155581014"></a>接口分类</p>
505</th>502</th>
506<th class="cellrowborder" valign="top" width="44.519999999999996%" id="mcps1.2.4.1.2"><p id="p12351955181013"><a name="p12351955181013"></a><a name="p12351955181013"></a>接口名称</p>503<th class="cellrowborder" valign="top" width="44.519999999999996%" id="mcps1.2.4.1.2"><p id="p12351955181013"><a name="p12351955181013"></a><a name="p12351955181013"></a>接口名称</p>
@@ -714,9 +711,8 @@
714<!-- end id3 -->711<!-- end id3 -->
715 712 
716<!-- npu="950" id4 -->713<!-- npu="950" id4 -->
717-**表4** 针对Ascend 950PR&950DT系列产品,支持的API范围714+**表4** 针对Ascend 950PR&950DT系列产品,支持的API范围<a name="table1798673512413"></a>
718 715 
719-<a name="table1798673512413"></a>
720<table><thead align="left"><tr id="row109864353413"><th class="cellrowborder" valign="top" width="29.630000000000003%" id="mcps1.2.4.1.1"><p id="p1098613524117"><a name="p1098613524117"></a><a name="p1098613524117"></a>接口分类</p>716<table><thead align="left"><tr id="row109864353413"><th class="cellrowborder" valign="top" width="29.630000000000003%" id="mcps1.2.4.1.1"><p id="p1098613524117"><a name="p1098613524117"></a><a name="p1098613524117"></a>接口分类</p>
721</th>717</th>
722<th class="cellrowborder" valign="top" width="44.519999999999996%" id="mcps1.2.4.1.2"><p id="p29871535174117"><a name="p29871535174117"></a><a name="p29871535174117"></a>接口名称</p>718<th class="cellrowborder" valign="top" width="44.519999999999996%" id="mcps1.2.4.1.2"><p id="p29871535174117"><a name="p29871535174117"></a><a name="p29871535174117"></a>接口名称</p>
@@ -1,8 +1,7 @@
1# 术语表<a name="ZH-CN_TOPIC_0000002191896252"></a>1# 术语表<a name="ZH-CN_TOPIC_0000002191896252"></a>
2 2 
3-**表1** 术语表3+**表1** 术语表<a name="table276611752817"></a>
4 4 
5-<a name="table276611752817"></a>
6<table><thead align="left"><tr id="row117661277289"><th class="cellrowborder" valign="top" width="22.978822978822976%" id="mcps1.2.3.1.1"><p id="p17766117132820"><a name="p17766117132820"></a><a name="p17766117132820"></a>术语/缩略语</p>5<table><thead align="left"><tr id="row117661277289"><th class="cellrowborder" valign="top" width="22.978822978822976%" id="mcps1.2.3.1.1"><p id="p17766117132820"><a name="p17766117132820"></a><a name="p17766117132820"></a>术语/缩略语</p>
7</th>6</th>
8<th class="cellrowborder" valign="top" width="77.02117702117702%" id="mcps1.2.3.1.2"><p id="p47664714282"><a name="p47664714282"></a><a name="p47664714282"></a>含义</p>7<th class="cellrowborder" valign="top" width="77.02117702117702%" id="mcps1.2.3.1.2"><p id="p47664714282"><a name="p47664714282"></a><a name="p47664714282"></a>含义</p>
@@ -43,8 +43,7 @@
43 43 
44![](../../../figures/zn_abc.png)44![](../../../figures/zn_abc.png)
45 45 
46-- **FRACTAL\_NZ/NZ**46+- **FRACTAL\_NZ/NZ**<a name="li19960204116136"></a>
47-<a name="li19960204116136"></a>
48 47 
49 FRACTAL\_NZ格式,简称NZ格式,是对一个Tensor最低两维(一个Tensor的所有维度,右侧为低维,左侧为高维)进行填充(pad)、拆分(reshape)和转置(transpose)操作后得到的格式。具体的转换过程如下:48 FRACTAL\_NZ格式,简称NZ格式,是对一个Tensor最低两维(一个Tensor的所有维度,右侧为低维,左侧为高维)进行填充(pad)、拆分(reshape)和转置(transpose)操作后得到的格式。具体的转换过程如下:
50 49 
@@ -19,9 +19,8 @@
19 19 
20Tensor是算子计算数据的容器,包含如下属性信息。20Tensor是算子计算数据的容器,包含如下属性信息。
21 21 
22-**表1** Tensor属性信息22+**表1** Tensor属性信息<a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_table196765568157"></a>
23 23 
24-<a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_table196765568157"></a>
25<table><thead align="left"><tr id="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_ra39c1f8353b1438b9aeba38c06570057"><th class="cellrowborder" valign="top" width="28.199999999999996%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_zh-cn_topic_0116955499_p6480165326"><a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_zh-cn_topic_0116955499_p6480165326"></a><a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_zh-cn_topic_0116955499_p6480165326"></a>属性</p>24<table><thead align="left"><tr id="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_ra39c1f8353b1438b9aeba38c06570057"><th class="cellrowborder" valign="top" width="28.199999999999996%" id="mcps1.2.3.1.1"><p id="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_zh-cn_topic_0116955499_p6480165326"><a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_zh-cn_topic_0116955499_p6480165326"></a><a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_zh-cn_topic_0116955499_p6480165326"></a>属性</p>
26</th>25</th>
27<th class="cellrowborder" valign="top" width="71.8%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_ab4f79bcaae2544bfad2388951768e902"><a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_ab4f79bcaae2544bfad2388951768e902"></a><a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_ab4f79bcaae2544bfad2388951768e902"></a>定义</p>26<th class="cellrowborder" valign="top" width="71.8%" id="mcps1.2.3.1.2"><p id="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_ab4f79bcaae2544bfad2388951768e902"><a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_ab4f79bcaae2544bfad2388951768e902"></a><a name="zh-cn_topic_0000001265237154_zh-cn_topic_0228422146_zh-cn_topic_0187054064_zh-cn_topic_0160787607_ab4f79bcaae2544bfad2388951768e902"></a>定义</p>
@@ -56,9 +55,8 @@ Tensor是算子计算数据的容器,包含如下属性信息。
56 55 
57形状的第一个元素对应张量最外层中括号中的元素个数,形状的第二个元素对应张量中从左边开始数第二个中括号中的元素个数,依此类推。例如:56形状的第一个元素对应张量最外层中括号中的元素个数,形状的第二个元素对应张量中从左边开始数第二个中括号中的元素个数,依此类推。例如:
58 57 
59-**表2** 张量的形状举例58+**表2** 张量的形状举例<a name="table08211915203"></a>
60 59 
61-<a name="table08211915203"></a>
62<table><thead align="left"><tr id="row482209142014"><th class="cellrowborder" valign="top" width="27.339999999999996%" id="mcps1.2.4.1.1"><p id="p28223992018"><a name="p28223992018"></a><a name="p28223992018"></a>张量</p>60<table><thead align="left"><tr id="row482209142014"><th class="cellrowborder" valign="top" width="27.339999999999996%" id="mcps1.2.4.1.1"><p id="p28223992018"><a name="p28223992018"></a><a name="p28223992018"></a>张量</p>
63</th>61</th>
64<th class="cellrowborder" valign="top" width="34.4%" id="mcps1.2.4.1.2"><p id="p1082213919205"><a name="p1082213919205"></a><a name="p1082213919205"></a>形状</p>62<th class="cellrowborder" valign="top" width="34.4%" id="mcps1.2.4.1.2"><p id="p1082213919205"><a name="p1082213919205"></a><a name="p1082213919205"></a>形状</p>