已合并
docs: 锚点与相邻文字合并至同一行 #6592
ai_xin创建于 28 天前
docs: 锚点与相邻文字合并至同一行 #6592
已合并
共 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** Hccl | 359 | +**表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 |  | 8 |  |
| 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的占比明显下降。 | 112 | - 优化后的Cube流水图如下所示,不使用随路ND2NZ指令对B矩阵进行格式转换后,MTE2的占比明显下降。 |
| 114 | - 优化后的Profiling数据如下,可以看到同时使用Cube单元和Vector单元,aic\_time最大耗时90.95us,其中aic\_mte2\_ratio占比明显降低。 | 113 | - 优化后的Profiling数据如下,可以看到同时使用Cube单元和Vector单元,aic\_time最大耗时90.95us,其中aic\_mte2\_ratio占比明显降低。 |
| 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> |
| @@ -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> |
Mdocs/zh/guide/operator_practice/best_practices/matmul_tuning_cases/matmul_multi_core_k_stagger.md+1-2
| @@ -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> |
Mdocs/zh/guide/operator_practice/best_practices/matmul_tuning_cases/matmul_multi_core_split_k.md+1-2
| @@ -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> |
Mdocs/zh/guide/operator_practice/best_practices/matmul_tuning_cases/matmul_optimization_overview.md+5-10
| @@ -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\_LEAKYRELU | 29 | - 算子类型(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 | ||
| 20 | MX格式的数据类型包含多种,例如MXFP8、MXFP4、MXFP16、MXINT4等。下表列举了[MxMatmul场景](#zh-cn_topic_0000002270097206_section310824820358)(全称Microscaling Matmul)支持的数据类型。 | 20 | MX格式的数据类型包含多种,例如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 |  | 8 |  |
| 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 | + |
| 12 | 1. **创建Tiling对象**。 | 12 | 1. **创建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 | ||
| 64 | 1. **创建Matmul对象。** | 63 | 1. **创建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 |  | 171 |  |
| 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):Add | 48 | - 算子类型(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):Add | 9 | - 算子类型(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> |
Mdocs/zh/guide/operator_practice/simd_operator_optimization/memory_access/l2_cache_mode_config.md+1-2
| @@ -38,9 +38,8 @@ L2 Cache的带宽相比GM的带宽有数倍的提升,因此当数据命中L2 C | |||
| 38 | 38 | ||
| 39 | Ascend C基于GlobalTensor提供了SetL2CacheHint接口,用户可以根据需要指定CacheMode。 | 39 | Ascend 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> |
Mdocs/zh/guide/operator_practice/simd_operator_optimization/vector_compute/low_latency_instruction.md+2-4
| @@ -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 | ||
| 31 | 3. 明确函数名称和参数。 | 31 | 3. 明确函数名称和参数。 |
| 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):Gather | 114 | - 算子类型(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) | |||
| 147 | export ASCENDC_BUILD_LOG_DIR=/home/build_log/ | 147 | export 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 | ||
| 3 | SIMD与SIMT混合编程涉及多种类型的函数,它们之间遵循严格的调用关系和层级约束。本章节将介绍整体函数类型和调用关系,再详细展开每种函数的定义、调用语法及约束。 | 3 | SIMD与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, | |||
| 56 | kernel_name<<<block_num, dyn_ub_size, stream>>>(args...); | 52 | kernel_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 | > |
Mdocs/zh/guide/programming_guide/advanced_programming/hardware_implementation/basic_architecture.md+2-4
| @@ -123,9 +123,8 @@ AI处理器中的计算资源要想发挥强劲算力,必要条件是保证输 | |||
| 123 | **图6** 存储单元<a name="fig2032418113610"></a> | 123 | **图6** 存储单元<a name="fig2032418113610"></a> |
| 124 |  | 124 |  |
| 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_ | |||
| 386 | bash run.sh -r npu -v <soc_version> -i <install_path> -b Debug -p <install-prefix> | 384 | bash 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 | ||
| 227 | SIMT编程提供了一系列运算符,用于执行数学运算。以下是支持的运算符列表。 | 217 | SIMT编程提供了一系列运算符,用于执行数学运算。以下是支持的运算符列表。 |
| 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 | ||
| 3 | Mask用于控制矢量计算中参与计算的元素个数,支持以下工作模式及配置方式: | 3 | Mask用于控制矢量计算中参与计算的元素个数,支持以下工作模式及配置方式: |
| 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 | ||
| 53 | Mask操作的使用方式如下: | 51 | Mask操作的使用方式如下: |
| 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 | ||
| 19 | AI Core中的计算单元主要包括Scalar、Vector和Cube三类。 | 19 | AI 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> |
Mdocs/zh/guide/technical_appendix/concepts_and_terms/neural_networks_and_operators/data_layout.md+1-2
| @@ -43,8 +43,7 @@ | |||
| 43 | 43 | ||
| 44 |  | 44 |  |
| 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 | ||
Mdocs/zh/guide/technical_appendix/concepts_and_terms/neural_networks_and_operators/operator_basics.md+2-4
| @@ -19,9 +19,8 @@ | |||
| 19 | 19 | ||
| 20 | Tensor是算子计算数据的容器,包含如下属性信息。 | 20 | Tensor是算子计算数据的容器,包含如下属性信息。 |
| 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> |