Star1069
1153
代码
代码
Issues113
Pull Requests312
Actions
项目讨论
Wiki
分析
Star1069
1153
已合并
docs: add blank lines around headings in docs/zh to meet format rules #6515
ai_xin创建于 19 天前
docs: add blank lines around headings in docs/zh to meet format rules #6515
已合并
ai_xin创建于 19 天前
从ai_xin/asc-devkitmaster合入到cann/asc-devkitmaster
共 107 个文件变更+189-0
Mdocs/zh/api/SIMD-API/adv_api/HCCL_communication/HCCL_Tiling/SetReduceType.md+1-0
@@ -26,6 +26,7 @@ uint32_t SetReduceType(uint32_t reduceType, uint8_t dstDataType = 0, uint8_t src
26- 非0表示设置失败。26- 非0表示设置失败。
27 27 
28## 约束说明28## 约束说明
29+ 
29<!-- npu="950" id3 -->30<!-- npu="950" id3 -->
30对于Ascend 950PR&950DT系列产品，AI CPU作为服务端时，该接口暂不支持使用。31对于Ascend 950PR&950DT系列产品，AI CPU作为服务端时，该接口暂不支持使用。
31<!-- end id3 -->32<!-- end id3 -->
Mdocs/zh/api/SIMD-API/adv_api/cube_compute/Matmul_Tiling/SetMadType.md+1-0
@@ -39,6 +39,7 @@ enum class MatrixMadType : int32_t {
39无39无
40 40 
41<!-- npu="950" id3 -->41<!-- npu="950" id3 -->
42+ 
42## 调用示例43## 调用示例
43 44 
44```45```
Mdocs/zh/api/SIMD-API/basic_api/cube_compute_ISASI/cube_compute_fractal_intro/aux_cube_fractal_format_details.md+2-0
@@ -1,6 +1,7 @@
1# 辅助矩阵分形格式详解<a name="ZH-CN_TOPIC_0000002538231118"></a>1# 辅助矩阵分形格式详解<a name="ZH-CN_TOPIC_0000002538231118"></a>
2 2 
3<!-- npu="950" id1 -->3<!-- npu="950" id1 -->
4+ 
4## MX系数矩阵<a name="zh-cn_topic_0000002563445163_section648585413300"></a>5## MX系数矩阵<a name="zh-cn_topic_0000002563445163_section648585413300"></a>
5 6 
6针对Ascend 950PR&950DT系列产品产品：7针对Ascend 950PR&950DT系列产品产品：
@@ -71,6 +72,7 @@ MXMmad为带有量化系数X的矩阵乘法，即左矩阵和右矩阵均有对
71<!-- end id1 -->72<!-- end id1 -->
72 73 
73<!-- npu="A3,910b" id2 -->74<!-- npu="A3,910b" id2 -->
75+ 
74## 4选2稀疏索引矩阵<a name="zh-cn_topic_0000002563445163_section1873692415233"></a>76## 4选2稀疏索引矩阵<a name="zh-cn_topic_0000002563445163_section1873692415233"></a>
75 77 
76针对Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品：78针对Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品：
Mdocs/zh/api/SIMD-API/basic_api/cube_compute_ISASI/cube_compute_fractal_intro/key_fractal_format_details.md+1-0
@@ -76,6 +76,7 @@
76<!-- end id20 -->76<!-- end id20 -->
77 77 
78<!-- npu="910b,A3,310b,310p,910,x90,9030" id10 -->78<!-- npu="910b,A3,310b,310p,910,x90,9030" id10 -->
79+ 
79## Zz格式（FRACTAL_Zz）<a name="zh-cn_topic_0000002545089965_section1798516571585"></a>80## Zz格式（FRACTAL_Zz）<a name="zh-cn_topic_0000002545089965_section1798516571585"></a>
80 81 
81- **定义：** 大Z外部row major（行主序）+ 小z内部row major（行主序），见[图4](#zh-cn_topic_0000002545089965_fig18376432205510)。82- **定义：** 大Z外部row major（行主序）+ 小z内部row major（行主序），见[图4](#zh-cn_topic_0000002545089965_fig18376432205510)。
Mdocs/zh/api/SIMD-API/basic_api/cube_compute_ISASI/cube_compute_load/LoadDataWithTranspose.md+2-0
@@ -156,6 +156,7 @@ Kirin X90处理器系列产品，支持数据类型：int8_t、half。
156<!-- npu="950" id19 -->156<!-- npu="950" id19 -->
157- 针对Ascend 950PR&950DT系列产品，推荐使用LoadData2dTransposeParamsV2作为参数，该参数具有更精细的搬运粒度。157- 针对Ascend 950PR&950DT系列产品，推荐使用LoadData2dTransposeParamsV2作为参数，该参数具有更精细的搬运粒度。
158<!-- end id19 -->158<!-- end id19 -->
159+ 
159## 关键特性说明<a id="zh-cn_topic_0000002543851571_section1891111310132"></a>160## 关键特性说明<a id="zh-cn_topic_0000002543851571_section1891111310132"></a>
160 161 
161### 功能和参数讲解162### 功能和参数讲解
@@ -514,6 +515,7 @@ for (int i = 0; i < CeilDivision(k, fractalShape[0]); ++i) {
514```515```
515 516 
516<!-- npu="950" id20 -->517<!-- npu="950" id20 -->
518+ 
517### 使用LoadData2dTransposeParamsV2结构体作为参数的场景示例519### 使用LoadData2dTransposeParamsV2结构体作为参数的场景示例
518 520 
519该示例使用了LoadData2dTransposeParamsV2结构体作为参数，输入a矩阵为int8_t类型，shape为[128,128]，输入数据格式为NZ，输入b矩阵为int8_t类型，shape为[128,256]，输入数据格式为NZ，输出c的类型为float。a矩阵从L1 Buffer（A1）搬运到L0A Buffer（A2）且不转置，b矩阵从L1 Buffer（B1）搬运到L0B Buffer（B2）并转置，示例仅展示接口调用过程，其余计算和搬运不作参考。完整示例请参考：[load_data_l12l0样例](../../../../../../../examples/01_simd_cpp_api/03_basic_api/03_matrix_compute/load_data_l12l0)。521该示例使用了LoadData2dTransposeParamsV2结构体作为参数，输入a矩阵为int8_t类型，shape为[128,128]，输入数据格式为NZ，输入b矩阵为int8_t类型，shape为[128,256]，输入数据格式为NZ，输出c的类型为float。a矩阵从L1 Buffer（A1）搬运到L0A Buffer（A2）且不转置，b矩阵从L1 Buffer（B1）搬运到L0B Buffer（B2）并转置，示例仅展示接口调用过程，其余计算和搬运不作参考。完整示例请参考：[load_data_l12l0样例](../../../../../../../examples/01_simd_cpp_api/03_basic_api/03_matrix_compute/load_data_l12l0)。
Mdocs/zh/api/SIMD-API/basic_api/cube_compute_ISASI/cube_compute_load/cube_compute_input_move_constraint.md+2-0
@@ -32,6 +32,7 @@ LoadDataWithSparse主要用于搬运以512Byte为单位存放的稠密权重矩
32矩阵计算过程中，常用分形支持总结如下表：32矩阵计算过程中，常用分形支持总结如下表：
33 33
34<!-- npu="910b,A3" id2 -->34<!-- npu="910b,A3" id2 -->
35+ 
35### 针对Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品36### 针对Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品
36 37 
37常用分形支持情况如下，供开发者参考：38常用分形支持情况如下，供开发者参考：
@@ -119,6 +120,7 @@ LoadDataWithSparse主要用于搬运以512Byte为单位存放的稠密权重矩
119<!-- end id2 -->120<!-- end id2 -->
120 121 
121<!-- npu="950" id3 -->122<!-- npu="950" id3 -->
123+ 
122### 针对Ascend 950PR&950DT系列产品124### 针对Ascend 950PR&950DT系列产品
123 125 
124DataCopy（GM-\>L1 Buffer）支持使能随路进行DN到NZ转换，一定会改变分形排布。126DataCopy（GM-\>L1 Buffer）支持使能随路进行DN到NZ转换，一定会改变分形排布。
Mdocs/zh/api/SIMD-API/basic_api/cube_compute_ISASI/cube_compute_store/DataCopy_L0CToL1.md+1-0
@@ -203,6 +203,7 @@ DataCopy矩阵搬出接口支持多种随路能力的组合，需要设置不同
203<!-- end id24 -->203<!-- end id24 -->
204 204 
205<!-- npu="950,A3,910b" id25 -->205<!-- npu="950,A3,910b" id25 -->
206+ 
206## 调用示例<a id="zh-cn_topic_0000002511188540_section088124295117"></a>207## 调用示例<a id="zh-cn_topic_0000002511188540_section088124295117"></a>
207 208 
208示例：通路L0C Buffer->L1 Buffer。输入A矩阵和B矩阵的数据类型为half，输出C矩阵为int8_t，不开启NZ2ND的格式转换，开启scalar量化。209示例：通路L0C Buffer->L1 Buffer。输入A矩阵和B矩阵的数据类型为half，输出C矩阵为int8_t，不开启NZ2ND的格式转换，开启scalar量化。
Mdocs/zh/api/SIMD-API/basic_api/cube_compute_ISASI/cube_compute_store/Fixpipe_L0CToUB.md+1-0
@@ -203,6 +203,7 @@ L0C Buffer到UB数据搬运提供矩阵搬出的组合接口Fixpipe，接口内
203<!-- end id18 -->203<!-- end id18 -->
204 204 
205<!-- npu="950" id20 -->205<!-- npu="950" id20 -->
206+ 
206## 关键特性说明207## 关键特性说明
207 208 
208以下图片和说明仅适用于Ascend 950PR&950DT系列产品：209以下图片和说明仅适用于Ascend 950PR&950DT系列产品：
Mdocs/zh/api/SIMD-API/basic_api/data_move_guide/GM_L1_or_L0_data_move.md+3-0
@@ -110,6 +110,7 @@ src和dst分别为源操作数和目的操作数；intriParams为ND2NZ搬运配
110<!-- end id1 -->110<!-- end id1 -->
111 111 
112<!-- npu="950" id2 -->112<!-- npu="950" id2 -->
113+ 
113## DataCopy（GMToL1随路转换-DN2NZ搬运）<a name="ZH-CN_TOPIC_0000002573902899"></a>114## DataCopy（GMToL1随路转换-DN2NZ搬运）<a name="ZH-CN_TOPIC_0000002573902899"></a>
114 115 
115该接口主要实现将数据从Global Memory搬运至L1 Buffer，并支持在数据搬运时进行[DN](../cube_compute_ISASI/cube_compute_fractal_intro/key_fractal_format_details.md)到[NZ](../cube_compute_ISASI/cube_compute_fractal_intro/key_fractal_format_details.md)格式的转换。116该接口主要实现将数据从Global Memory搬运至L1 Buffer，并支持在数据搬运时进行[DN](../cube_compute_ISASI/cube_compute_fractal_intro/key_fractal_format_details.md)到[NZ](../cube_compute_ISASI/cube_compute_fractal_intro/key_fractal_format_details.md)格式的转换。
@@ -132,6 +133,7 @@ src和dst分别为源操作数和目的操作数；intriParams为DN2NZ搬运配
132<!-- end id2 -->133<!-- end id2 -->
133 134 
134<!-- npu="950" id3 -->135<!-- npu="950" id3 -->
136+ 
135## DataCopyPad（GMToL1非对齐数据搬运）<a name="ZH-CN_TOPIC_0000002543262891"></a>137## DataCopyPad（GMToL1非对齐数据搬运）<a name="ZH-CN_TOPIC_0000002543262891"></a>
136 138 
137该接口主要实现将数据从Global Memory搬运至L1 Buffer，并支持在数据搬运时提供非对齐数据搬运功能，可以根据开发者的需要自行填充数据。139该接口主要实现将数据从Global Memory搬运至L1 Buffer，并支持在数据搬运时提供非对齐数据搬运功能，可以根据开发者的需要自行填充数据。
@@ -172,6 +174,7 @@ __aicore__ inline void LoadData(const LocalTensor<T>& dst, const GlobalTensor<T>
172```174```
173 175 
174<!-- npu="950" id4 -->176<!-- npu="950" id4 -->
177+ 
175## GMToL1 NZ数据搬运(LoadData（2D矩阵搬运V2）)178## GMToL1 NZ数据搬运(LoadData（2D矩阵搬运V2）)
176 179 
177该接口主要实现将数据从Global Memory搬运至L1 Buffer，负责完成普通矩阵计算所需的2D格式数据的搬运，以大小为512Byte的数据分形为单位进行搬运。180该接口主要实现将数据从Global Memory搬运至L1 Buffer，负责完成普通矩阵计算所需的2D格式数据的搬运，以大小为512Byte的数据分形为单位进行搬运。
Mdocs/zh/api/SIMD-API/basic_api/data_move_guide/L1_L0_data_move.md+7-0
@@ -43,6 +43,7 @@ __aicore__ inline void LoadData(const LocalTensor<T>& dst, const LocalTensor<T>&
43```43```
44 44 
45<!-- npu="950" id1 -->45<!-- npu="950" id1 -->
46+ 
46## L1ToL0A 2D格式分形矩阵搬运(LoadData（2D矩阵搬运V2）)47## L1ToL0A 2D格式分形矩阵搬运(LoadData（2D矩阵搬运V2）)
47 48 
48LoadData（2D矩阵搬运V2）能够实现L1 Buffer到L0A Buffer之间的数据搬运。负责完成普通矩阵计算所需的2D格式数据的搬运，以大小为512Byte的数据分形为单位进行搬运。49LoadData（2D矩阵搬运V2）能够实现L1 Buffer到L0A Buffer之间的数据搬运。负责完成普通矩阵计算所需的2D格式数据的搬运，以大小为512Byte的数据分形为单位进行搬运。
@@ -65,6 +66,7 @@ src和dst分别为源操作数和目的操作数；loadDataParams为搬运参数
65<!-- end id1 -->66<!-- end id1 -->
66 67 
67<!-- npu="950" id2 -->68<!-- npu="950" id2 -->
69+ 
68## L1ToL0A 2D格式分形矩阵搬运(LoadData（MX矩阵搬运）)70## L1ToL0A 2D格式分形矩阵搬运(LoadData（MX矩阵搬运）)
69 71 
70LoadData（MX矩阵搬运）负责完成矩阵计算所需的左右矩阵数据和对应的左右量化系数矩阵数据的搬运，其中左右矩阵数据以大小为512Byte的数据分形为单位进行搬运，左右量化系数矩阵以32Byte的数据分形为单位进行搬运。72LoadData（MX矩阵搬运）负责完成矩阵计算所需的左右矩阵数据和对应的左右量化系数矩阵数据的搬运，其中左右矩阵数据以大小为512Byte的数据分形为单位进行搬运，左右量化系数矩阵以32Byte的数据分形为单位进行搬运。
@@ -157,6 +159,7 @@ src和dst分别为源操作数和目的操作数；loadDataParams为搬运参数
157<!-- end id3 -->159<!-- end id3 -->
158 160 
159<!-- npu="950" id1 -->161<!-- npu="950" id1 -->
162+ 
160## L1ToL0A 3D格式分形矩阵搬运(LoadDataWithStride)<a name="ZH-CN_TOPIC_0000002573312819"></a>163## L1ToL0A 3D格式分形矩阵搬运(LoadDataWithStride)<a name="ZH-CN_TOPIC_0000002573312819"></a>
161 164 
162LoadDataWithStride用于将NC1HWC0格式的Feature Map完成Image to Column展开，然后再从展开后的二维矩阵中选取指定数据块搬入对应内存位置，支持配置输出矩阵K轴方向偏移量。165LoadDataWithStride用于将NC1HWC0格式的Feature Map完成Image to Column展开，然后再从展开后的二维矩阵中选取指定数据块搬入对应内存位置，支持配置输出矩阵K轴方向偏移量。
@@ -197,6 +200,7 @@ __aicore__ inline void LoadData(const LocalTensor<T>& dst, const LocalTensor<T>&
197```200```
198 201 
199<!-- npu="950" id4 -->202<!-- npu="950" id4 -->
203+ 
200## L1ToL0B 2D格式分形矩阵搬运(LoadData（2D矩阵搬运V2）)204## L1ToL0B 2D格式分形矩阵搬运(LoadData（2D矩阵搬运V2）)
201 205 
202LoadData（2D矩阵搬运V2）能够实现L1 Buffer到L0B Buffer之间的数据搬运。负责完成普通矩阵计算所需的2D格式数据的搬运，以大小为512Byte的数据分形为单位进行搬运。206LoadData（2D矩阵搬运V2）能够实现L1 Buffer到L0B Buffer之间的数据搬运。负责完成普通矩阵计算所需的2D格式数据的搬运，以大小为512Byte的数据分形为单位进行搬运。
@@ -219,6 +223,7 @@ src和dst分别为源操作数和目的操作数；loadDataParams为搬运参数
219<!-- end id4 -->223<!-- end id4 -->
220 224 
221<!-- npu="950" id5 -->225<!-- npu="950" id5 -->
226+ 
222## L1ToL0B 2D格式分形矩阵搬运(LoadData（MX矩阵搬运）)227## L1ToL0B 2D格式分形矩阵搬运(LoadData（MX矩阵搬运）)
223 228 
224LoadData（MX矩阵搬运）负责完成矩阵计算所需的左右矩阵数据和对应的左右量化系数矩阵数据的搬运，其中左右矩阵数据以大小为512Byte的数据分形为单位进行搬运，左右量化系数矩阵以32Byte的数据分形为单位进行搬运。229LoadData（MX矩阵搬运）负责完成矩阵计算所需的左右矩阵数据和对应的左右量化系数矩阵数据的搬运，其中左右矩阵数据以大小为512Byte的数据分形为单位进行搬运，左右量化系数矩阵以32Byte的数据分形为单位进行搬运。
@@ -311,6 +316,7 @@ src和dst分别为源操作数和目的操作数；loadDataParams为搬运参数
311<!-- end id6 -->316<!-- end id6 -->
312 317 
313<!-- npu="950" id12 -->318<!-- npu="950" id12 -->
319+ 
314## L1ToL0B 3D格式分形矩阵搬运(LoadDataWithStride)<a name="ZH-CN_TOPIC_0000002573312819"></a>320## L1ToL0B 3D格式分形矩阵搬运(LoadDataWithStride)<a name="ZH-CN_TOPIC_0000002573312819"></a>
315 321 
316LoadDataWithStride用于将NC1HWC0格式的Feature Map完成Image to Column展开，然后再从展开后的二维矩阵中选取指定数据块搬入对应内存位置，支持配置输出矩阵K轴方向偏移量。322LoadDataWithStride用于将NC1HWC0格式的Feature Map完成Image to Column展开，然后再从展开后的二维矩阵中选取指定数据块搬入对应内存位置，支持配置输出矩阵K轴方向偏移量。
@@ -333,6 +339,7 @@ src和dst分别为源操作数和目的操作数；loadDataParams为搬运参数
333<!-- end id12 -->339<!-- end id12 -->
334 340 
335<!-- npu="A3,910b" id9 -->341<!-- npu="A3,910b" id9 -->
342+ 
336## L1ToL0B 稠密权重矩阵搬运(LoadDataWithSparse)<a name="ZH-CN_TOPIC_0000002574022821"></a>343## L1ToL0B 稠密权重矩阵搬运(LoadDataWithSparse)<a name="ZH-CN_TOPIC_0000002574022821"></a>
337 344 
338用于从L1中搬运以512Byte为单位存放的稠密权重矩阵到L0B里，同时搬运以128Byte为单位的索引矩阵到内置的专用buffer空间（用于后续MmadWithSparse接口进行读取）。345用于从L1中搬运以512Byte为单位存放的稠密权重矩阵到L0B里，同时搬运以128Byte为单位的索引矩阵到内置的专用buffer空间（用于后续MmadWithSparse接口进行读取）。
Mdocs/zh/api/SIMD-API/basic_api/data_move_guide/L1_or_L0C_UB_data_move.md+3-0
@@ -92,6 +92,7 @@ __aicore__ inline void DataCopyPad(const LocalTensor<T>& dst, const LocalTensor<
92```92```
93 93 
94<!-- npu="310p" id1 -->94<!-- npu="310p" id1 -->
95+ 
95## DataCopy（L0CToUB随路量化激活搬运）96## DataCopy（L0CToUB随路量化激活搬运）
96 97 
97该接口主要实现将数据从L0C Buffer搬运至UB，并支持多种随路能力的组合，接口内包含了设置寄存器与数据搬运能力。98该接口主要实现将数据从L0C Buffer搬运至UB，并支持多种随路能力的组合，接口内包含了设置寄存器与数据搬运能力。
@@ -152,6 +153,7 @@ src和dst分别为源操作数和目的操作数；intriParams为搬运参数，
152 ```153 ```
153 154 
154<!-- npu="950" id2 -->155<!-- npu="950" id2 -->
156+ 
155## DataCopyL1ToUB（L1ToUB连续数据搬运）157## DataCopyL1ToUB（L1ToUB连续数据搬运）
156 158 
157DataCopyL1ToUB能够实现L1 Buffer到UB的连续数据搬运，数据搬运时格式和内容保持不变。159DataCopyL1ToUB能够实现L1 Buffer到UB的连续数据搬运，数据搬运时格式和内容保持不变。
@@ -174,6 +176,7 @@ __aicore__ inline void DataCopyL1ToUB(const LocalTensor<T>& dst, const LocalTens
174<!-- end id2 -->176<!-- end id2 -->
175 177 
176<!-- npu="950" id3 -->178<!-- npu="950" id3 -->
179+ 
177## DataCopyL1ToUB（L1ToUB高维切分数据搬运）180## DataCopyL1ToUB（L1ToUB高维切分数据搬运）
178 181 
179DataCopyL1ToUB能够实现L1 Buffer到UB的连续数据搬运和非连续数据搬运，数据搬运时格式和内容保持不变。182DataCopyL1ToUB能够实现L1 Buffer到UB的连续数据搬运和非连续数据搬运，数据搬运时格式和内容保持不变。
Mdocs/zh/api/SIMD-API/basic_api/data_structures/builtin_data_types.md+1-0
@@ -381,6 +381,7 @@ $$
381 | 0 | 8'hFE | 23'h7FFFFF | 最大正规格数 |381 | 0 | 8'hFE | 23'h7FFFFF | 最大正规格数 |
382 382 
383<!-- npu="950" id1 -->383<!-- npu="950" id1 -->
384+ 
384## 复数<a name="section8666633175111"></a>385## 复数<a name="section8666633175111"></a>
385 386 
386Ascend C提供了复数数据类型：complex32和complex64。387Ascend C提供了复数数据类型：complex32和complex64。
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/basic_arithmetic/Div.md+1-0
@@ -208,6 +208,7 @@ $dst_i = src0_i / src1_i$
208<!-- end id32 -->208<!-- end id32 -->
209 209 
210<!-- npu="950" id35 -->210<!-- npu="950" id35 -->
211+ 
211## 关键特性<a id="div-key-features"></a>212## 关键特性<a id="div-key-features"></a>
212 213 
213针对Ascend 950PR&950DT系列产品，有如下关键特性：214针对Ascend 950PR&950DT系列产品，有如下关键特性：
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/basic_arithmetic/Exp.md+1-0
@@ -178,6 +178,7 @@ T支持的数据类型为：half、float。
178<!-- end id21 -->178<!-- end id21 -->
179 179 
180<!-- npu="950" id24 -->180<!-- npu="950" id24 -->
181+ 
181## 关键特性<a id="exp-key-features"></a>182## 关键特性<a id="exp-key-features"></a>
182 183 
183针对Ascend 950PR&950DT系列产品，有如下关键特性：184针对Ascend 950PR&950DT系列产品，有如下关键特性：
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/basic_arithmetic/Ln.md+1-0
@@ -180,6 +180,7 @@ T支持的数据类型为：half、float。
180<!-- end id21 -->180<!-- end id21 -->
181 181 
182<!-- npu="950" id24 -->182<!-- npu="950" id24 -->
183+ 
183## 关键特性<a id="ln-key-features"></a>184## 关键特性<a id="ln-key-features"></a>
184 185 
185针对Ascend 950PR&950DT系列产品，有如下关键特性：186针对Ascend 950PR&950DT系列产品，有如下关键特性：
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/basic_arithmetic/Reciprocal.md+1-0
@@ -204,6 +204,7 @@ $dst_i = \frac{1}{src_i}$
204- 使用Reciprocal时，half的算子结果对比误差不满足双千分之一的要求，float的算子结果对比误差不满足双万分之一的要求，如果需要高精度，建议使用[Div](Div.md)替代实现。204- 使用Reciprocal时，half的算子结果对比误差不满足双千分之一的要求，float的算子结果对比误差不满足双万分之一的要求，如果需要高精度，建议使用[Div](Div.md)替代实现。
205 205 
206<!-- npu="950" id35 -->206<!-- npu="950" id35 -->
207+ 
207## 关键特性<a id="reciprocal-key-features"></a>208## 关键特性<a id="reciprocal-key-features"></a>
208 209 
209针对Ascend 950PR&950DT系列产品，有如下关键特性：210针对Ascend 950PR&950DT系列产品，有如下关键特性：
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/basic_arithmetic/Rsqrt.md+1-0
@@ -181,6 +181,7 @@ T支持的数据类型为：half、float。
181- 使用Rsqrt时，half的算子结果对比误差不满足双千分之一的要求，float的算子结果对比误差不满足双万分之一的要求，如果需要高精度，建议使用[Div](Div.md)和[Sqrt](Sqrt.md)替代实现。181- 使用Rsqrt时，half的算子结果对比误差不满足双千分之一的要求，float的算子结果对比误差不满足双万分之一的要求，如果需要高精度，建议使用[Div](Div.md)和[Sqrt](Sqrt.md)替代实现。
182 182 
183<!-- npu="950" id35 -->183<!-- npu="950" id35 -->
184+ 
184## 关键特性<a id="rsqrt-key-features"></a>185## 关键特性<a id="rsqrt-key-features"></a>
185 186 
186针对Ascend 950PR&950DT系列产品，有如下关键特性：187针对Ascend 950PR&950DT系列产品，有如下关键特性：
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/basic_arithmetic/Sqrt.md+1-0
@@ -181,6 +181,7 @@ T支持的数据类型为：half、float。
181- 如果src中的数值为非正数，可能会产生未知结果。181- 如果src中的数值为非正数，可能会产生未知结果。
182 182 
183<!-- npu="950" id35 -->183<!-- npu="950" id35 -->
184+ 
184## 关键特性<a id="sqrt-key-features"></a>185## 关键特性<a id="sqrt-key-features"></a>
185 186 
186针对Ascend 950PR&950DT系列产品，有如下关键特性：187针对Ascend 950PR&950DT系列产品，有如下关键特性：
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/composite_compute/AbsSub_ISASI.md+1-0
@@ -59,6 +59,7 @@ __aicore__ inline void AbsSub(const LocalTensor<T> &dst, const LocalTensor<T> &s
59| count | 输入 | 参与计算的元素个数。 |59| count | 输入 | 参与计算的元素个数。 |
60 60 
61## 数据类型61## 数据类型
62+ 
62<!-- npu="950" id8 -->63<!-- npu="950" id8 -->
63Ascend 950PR&950DT系列产品，T支持的数据类型为：half、float。64Ascend 950PR&950DT系列产品，T支持的数据类型为：half、float。
64<!-- end id8 -->65<!-- end id8 -->
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/composite_compute/ExpSub_ISASI.md+1-0
@@ -64,6 +64,7 @@ __aicore__ inline void ExpSub(const LocalTensor<T>& dst, const LocalTensor<U>& s
64| count | 输入 | 参与计算的元素个数。 |64| count | 输入 | 参与计算的元素个数。 |
65 65 
66## 数据类型66## 数据类型
67+ 
67<!-- npu="950" id8 -->68<!-- npu="950" id8 -->
68Ascend 950PR&950DT系列产品，T和U支持的数据类型为：half、float。69Ascend 950PR&950DT系列产品，T和U支持的数据类型为：half、float。
69<!-- end id8 -->70<!-- end id8 -->
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/composite_compute/MulsCast_ISASI.md+1-0
@@ -84,6 +84,7 @@ __aicore__ inline void MulsCast(const T2 &dst, const T3 &src0, const T4 &src1, c
84| float | half | 将源操作数按照CAST_ROUND模式取到half所能表示的数，以half格式（溢出默认按照饱和处理）存入dst中。 |84| float | half | 将源操作数按照CAST_ROUND模式取到half所能表示的数，以half格式（溢出默认按照饱和处理）存入dst中。 |
85 85 
86## 数据类型86## 数据类型
87+ 
87<!-- npu="950" id8 -->88<!-- npu="950" id8 -->
88Ascend 950PR&950DT系列产品，src支持的数据类型为：float。dst支持的数据类型为：half。89Ascend 950PR&950DT系列产品，src支持的数据类型为：float。dst支持的数据类型为：half。
89<!-- end id8 -->90<!-- end id8 -->
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/data_padding/Brcb.md+1-0
@@ -140,6 +140,7 @@ __aicore__ inline void Brcb(const LocalTensor<T>& dst, const LocalTensor<T>& src
140 - 针对Ascend 950PR&950DT系列产品，该接口通过VF调用[Reg矢量计算API](../../reg_vector_compute/reg_vector_compute.md)实现兼容，当参数repeatTime取值为0时，不保证该接口将被视为NOP（空操作）。140 - 针对Ascend 950PR&950DT系列产品，该接口通过VF调用[Reg矢量计算API](../../reg_vector_compute/reg_vector_compute.md)实现兼容，当参数repeatTime取值为0时，不保证该接口将被视为NOP（空操作）。
141 <!-- end id21 -->141 <!-- end id21 -->
142<!-- end id17 -->142<!-- end id17 -->
143+ 
143## 关键特性说明144## 关键特性说明
144 145 
145源操作数不支持Stride参数配置，始终采用连续读。146源操作数不支持Stride参数配置，始终采用连续读。
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/logical_compute/And.md+1-0
@@ -165,6 +165,7 @@ $dst_i = src0_i \& src1_i$
165 - 针对Ascend 950PR&950DT系列产品，该接口通过VF调用[Reg矢量计算API](../../reg_vector_compute/reg_vector_compute.md)实现兼容，当参数count或repeatTime取值为0时，不保证该接口将被视为NOP（空操作）。165 - 针对Ascend 950PR&950DT系列产品，该接口通过VF调用[Reg矢量计算API](../../reg_vector_compute/reg_vector_compute.md)实现兼容，当参数count或repeatTime取值为0时，不保证该接口将被视为NOP（空操作）。
166 <!-- end id30 -->166 <!-- end id30 -->
167<!-- end id26 -->167<!-- end id26 -->
168+ 
168## 调用示例<a name="section642mcpsimp"></a>169## 调用示例<a name="section642mcpsimp"></a>
169 170 
170- tensor高维切分计算样例-mask连续模式171- tensor高维切分计算样例-mask连续模式
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/logical_compute/ShiftLeft_shift_amount_tensor.md+1-0
@@ -58,6 +58,7 @@ __aicore__ inline void ShiftLeft(const LocalTensor<T>& dst, const LocalTensor<T>
58| count | 输入 | 参与计算的元素个数。 |58| count | 输入 | 参与计算的元素个数。 |
59 59 
60## 数据类型60## 数据类型
61+ 
61<!-- npu="950" id8 -->62<!-- npu="950" id8 -->
62Ascend 950PR&950DT系列产品，T支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。U支持的数据类型为：int8_t、int16_t、int32_t、int64_t。63Ascend 950PR&950DT系列产品，T支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。U支持的数据类型为：int8_t、int16_t、int32_t、int64_t。
63<!-- end id8 -->64<!-- end id8 -->
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/logical_compute/ShiftRight_shift_amount_tensor.md+1-0
@@ -68,6 +68,7 @@ __aicore__ inline void ShiftRight(const LocalTensor<T>& dst, const LocalTensor<T
68 68 
69 69 
70## 数据类型70## 数据类型
71+ 
71<!-- npu="950" id8 -->72<!-- npu="950" id8 -->
72Ascend 950PR&950DT系列产品，T支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。U支持的数据类型为：int8_t、int16_t、int32_t、int64_t。73Ascend 950PR&950DT系列产品，T支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t、int64_t、uint64_t。U支持的数据类型为：int8_t、int16_t、int32_t、int64_t。
73<!-- end id8 -->74<!-- end id8 -->
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/reduction_compute/ReduceMax.md+1-0
@@ -91,6 +91,7 @@
91**注：以上高维切分相关参数`mask`，`repeatTime`，`srcRepStride`请参考[高维切分](../SIMD_compute/high_dim_split.md)中的介绍。**91**注：以上高维切分相关参数`mask`，`repeatTime`，`srcRepStride`请参考[高维切分](../SIMD_compute/high_dim_split.md)中的介绍。**
92 92 
93## 数据类型93## 数据类型
94+ 
94支持的数据类型如下：95支持的数据类型如下：
95 96 
96<!-- npu="950" id10 -->97<!-- npu="950" id10 -->
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/scatter_gather/Gather.md+1-0
@@ -90,6 +90,7 @@ def Gather(dst, src, count, srcOffset, srcBaseAddr):
90| dstRepStride | 输入 | 相邻迭代间的地址步长，单位是DataBlock（32Bytes）。 |90| dstRepStride | 输入 | 相邻迭代间的地址步长，单位是DataBlock（32Bytes）。 |
91 91 
92## 数据类型92## 数据类型
93+ 
93<!-- npu="950" id8 -->94<!-- npu="950" id8 -->
94Ascend 950PR&950DT系列产品，支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float、int64_t、uint64_t。数据类型int8\_t、uint8\_t仅支持tensor前n个数据计算接口。95Ascend 950PR&950DT系列产品，支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float、int64_t、uint64_t。数据类型int8\_t、uint8\_t仅支持tensor前n个数据计算接口。
95<!-- end id8 -->96<!-- end id8 -->
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/scatter_gather/Gatherb_ISASI.md+1-0
@@ -149,6 +149,7 @@ Atlas 200I/500 A2推理产品，支持的数据类型为：int8_t、uint8_t、in
149 - Atlas A2系列产品149 - Atlas A2系列产品
150 <!-- end id20 -->150 <!-- end id20 -->
151<!-- end id18 -->151<!-- end id18 -->
152+ 
152## 调用示例<a name="section11276201527"></a>153## 调用示例<a name="section11276201527"></a>
153 154 
154完整使用样例请参见[Gather类样例](../../../../../../../examples/01_simd_cpp_api/03_basic_api/01_memory_vector_compute/gather)场景四。155完整使用样例请参见[Gather类样例](../../../../../../../examples/01_simd_cpp_api/03_basic_api/01_memory_vector_compute/gather)场景四。
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/type_conversion/Truncate_ISASI.md+1-0
@@ -57,6 +57,7 @@ __aicore__ inline void Truncate(const LocalTensor<T> &dst, const LocalTensor<T>
57| count | 输入 | 参与计算的元素个数。 |57| count | 输入 | 参与计算的元素个数。 |
58 58 
59## 数据类型59## 数据类型
60+ 
60<!-- npu="950" id8 -->61<!-- npu="950" id8 -->
61- 针对Ascend 950PR&950DT系列产品，T支持的数据类型为：half、float、bfloat16_t。62- 针对Ascend 950PR&950DT系列产品，T支持的数据类型为：half、float、bfloat16_t。
62<!-- end id8 -->63<!-- end id8 -->
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/basic_arithmetic/Mins.md+1-0
@@ -63,6 +63,7 @@ __simd_callee__ inline void Mins(S& dstReg, S& srcReg, U scalarValue, MaskReg& m
63| mask | 输入 | 源操作数元素操作的有效指示，详细说明请参考[MaskReg](../register_data_types/MaskReg.md)。 |63| mask | 输入 | 源操作数元素操作的有效指示，详细说明请参考[MaskReg](../register_data_types/MaskReg.md)。 |
64 64 
65## 数据类型65## 数据类型
66+ 
66目的操作数与源操作数的数据类型需要保持一致。支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float、int64_t、uint64_t。67目的操作数与源操作数的数据类型需要保持一致。支持的数据类型为：int8_t、uint8_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float、int64_t、uint64_t。
67 68 
68## 返回值说明<a name="section640mcpsimp"></a>69## 返回值说明<a name="section640mcpsimp"></a>
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/key_features/key_features.md+4-0
@@ -5,6 +5,7 @@
5Reg矢量计算架构中，Vector Function（VF）是实现高性能向量计算的核心载体。VF函数中可以包含最多四层嵌套循环，每层循环中还可以包含多个串行循环。VF循环对控制结构的支持有限，仅支持for循环和条件判断，不支持switch、do-while和while-do等其他控制结构。当VF函数中的循环满足Hardware Loop编码规范会被编译器优化为Hardware Loop，提升整体的编码性能，否则它的循环逻辑会由迭代变量和条件判断语句构成Software Loop，无法开启VF循环优化。5Reg矢量计算架构中，Vector Function（VF）是实现高性能向量计算的核心载体。VF函数中可以包含最多四层嵌套循环，每层循环中还可以包含多个串行循环。VF循环对控制结构的支持有限，仅支持for循环和条件判断，不支持switch、do-while和while-do等其他控制结构。当VF函数中的循环满足Hardware Loop编码规范会被编译器优化为Hardware Loop，提升整体的编码性能，否则它的循环逻辑会由迭代变量和条件判断语句构成Software Loop，无法开启VF循环优化。
6 6 
7### Hardware Loop编码规范7### Hardware Loop编码规范
8+ 
8- Hardware Loop最多能支持的嵌套层数为4层；支持串行的Hardware Loop循环。9- Hardware Loop最多能支持的嵌套层数为4层；支持串行的Hardware Loop循环。
9- 迭代变量类型<br>VF内所有Loop的迭代变量必须是uint16_t类型。10- 迭代变量类型<br>VF内所有Loop的迭代变量必须是uint16_t类型。
10- 起始值与步长<br>循环起始值从0开始。<br>每次迭代的步长必须是递增1。11- 起始值与步长<br>循环起始值从0开始。<br>每次迭代的步长必须是递增1。
@@ -13,6 +14,7 @@ Reg矢量计算架构中，Vector Function（VF）是实现高性能向量计算
13- 若要利用外层循环的计数作为循环边界，将外层循环计数器赋值给标量后作为内层循环边界。14- 若要利用外层循环的计数作为循环边界，将外层循环计数器赋值给标量后作为内层循环边界。
14 15 
15## 指令双发优化16## 指令双发优化
17+ 
16Reg矢量计算接口指令单双发的情况请参考[Reg矢量计算接口指令单双发汇总](../../../../appendix/reg_vector_compute_interface_instruction_single_dual_issue_summary.md)。18Reg矢量计算接口指令单双发的情况请参考[Reg矢量计算接口指令单双发汇总](../../../../appendix/reg_vector_compute_interface_instruction_single_dual_issue_summary.md)。
17 19 
18指令双发优化的性能优化详细内容请参考[算子实践参考-指令双发优化](../../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/dual_issue_optimization.md)。20指令双发优化的性能优化详细内容请参考[算子实践参考-指令双发优化](../../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/dual_issue_optimization.md)。
@@ -32,11 +34,13 @@ VF并不是写的越长，把所有运算都放在一个for循环内就好，需
32如果循环内存在依赖关系过多的指令，指令无法并发执行，即无法启用双发特性。可以通过展开循环，提升指令双发能力，贴近硬件乱序执行的特性；减少指令因为寄存器资源未到位而产生的等待。34如果循环内存在依赖关系过多的指令，指令无法并发执行，即无法启用双发特性。可以通过展开循环，提升指令双发能力，贴近硬件乱序执行的特性；减少指令因为寄存器资源未到位而产生的等待。
33 35 
34## VF融合36## VF融合
37+ 
35VF融合是将代码中多个VF函数融合成一个VF函数，有效提升性能。VF融合特性是编译器优化特性，通过编译选项--cce-simd-vf-fusion=true开启，VF自动融合会借助Loop Fuse算法，将VF转换成Loop形态，然后将控制流等价（Control-Flow-Equivalent）的VF进行融合，最后将VF进行还原。编译器首先会做融合前的合法性检查，判断两个VF是否等价，Main侧中间代码是否能在VF内执行以及融合后是否可产生正收益。如果满足VF融合条件，编译器会自动执行VF融合优化，为保证融合后的VF执行逻辑与语义与融合前一致，会在原来两个VF之间保守地插入同步指令，编译器还会尝试外提、合并融合后的VF中的指令，对VF代码进行优化。融合策略是能融尽融，用户按照符合融合的合法性检查的模式进行编码，可以增加VF融合的机会。用户也可以参考融合原理手动进行融合优化。38VF融合是将代码中多个VF函数融合成一个VF函数，有效提升性能。VF融合特性是编译器优化特性，通过编译选项--cce-simd-vf-fusion=true开启，VF自动融合会借助Loop Fuse算法，将VF转换成Loop形态，然后将控制流等价（Control-Flow-Equivalent）的VF进行融合，最后将VF进行还原。编译器首先会做融合前的合法性检查，判断两个VF是否等价，Main侧中间代码是否能在VF内执行以及融合后是否可产生正收益。如果满足VF融合条件，编译器会自动执行VF融合优化，为保证融合后的VF执行逻辑与语义与融合前一致，会在原来两个VF之间保守地插入同步指令，编译器还会尝试外提、合并融合后的VF中的指令，对VF代码进行优化。融合策略是能融尽融，用户按照符合融合的合法性检查的模式进行编码，可以增加VF融合的机会。用户也可以参考融合原理手动进行融合优化。
36 39 
37VF融合的性能优化详细内容请参考[算子实践参考-VF融合优化](../../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/vf_fusion_optimization.md)。40VF融合的性能优化详细内容请参考[算子实践参考-VF融合优化](../../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/vf_fusion_optimization.md)。
38 41 
39## 数据搬运优化42## 数据搬运优化
43+ 
40Reg矢量计算API提供了下表所示多种搬运指令，合理选择搬运接口以及利用接口能力进行优化。连续非对齐场景优化的性能优化详细内容请参考[算子实践参考-连续非对齐场景优化](../../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/continuous_unaligned_optimization.md)。44Reg矢量计算API提供了下表所示多种搬运指令，合理选择搬运接口以及利用接口能力进行优化。连续非对齐场景优化的性能优化详细内容请参考[算子实践参考-连续非对齐场景优化](../../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/continuous_unaligned_optimization.md)。
41 45 
42| 场景 | 描述 |46| 场景 | 描述 |
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/overview.md+2-0
@@ -1,7 +1,9 @@
1# 概述1# 概述
2+ 
2Reg矢量计算API是面向Reg矢量计算架构开发的API，用户可以通过该API直接对芯片中涉及Vector计算的寄存器进行操作，实现更大的灵活性和更好的性能。Reg矢量计算API与Memory矢量计算API功能相似，但与Memory矢量计算API输入和输出数据必须为LocalTensor不同，Reg矢量计算API的输入或输出数据均为Reg矢量计算寄存器。对于计算类API，其功能是从给定的寄存器获取数据，进行计算，并将结果保存在给定的寄存器。对于搬运类API，其功能是实现Unified Buffer（UB）和寄存器的数据搬运。由此可见，Reg矢量计算API相较于Memory矢量计算API，将数据搬运和Reg计算过程交给用户自主控制，从而实现更大的开发自由度。3Reg矢量计算API是面向Reg矢量计算架构开发的API，用户可以通过该API直接对芯片中涉及Vector计算的寄存器进行操作，实现更大的灵活性和更好的性能。Reg矢量计算API与Memory矢量计算API功能相似，但与Memory矢量计算API输入和输出数据必须为LocalTensor不同，Reg矢量计算API的输入或输出数据均为Reg矢量计算寄存器。对于计算类API，其功能是从给定的寄存器获取数据，进行计算，并将结果保存在给定的寄存器。对于搬运类API，其功能是实现Unified Buffer（UB）和寄存器的数据搬运。由此可见，Reg矢量计算API相较于Memory矢量计算API，将数据搬运和Reg计算过程交给用户自主控制，从而实现更大的开发自由度。
3 4 
4## 如何使用Reg矢量计算API5## 如何使用Reg矢量计算API
6+ 
5基于寄存器的编程模型是指将数据从LocalTensor通过Reg搬运指令（Reg数据搬入）加载到寄存器中，进行复杂的数学计算Compute后通过Reg搬运指令（Reg数据搬出）搬出到LocalTensor中，所有的计算逻辑均在寄存器中完成，从而减少LocalTensor间的数据搬运，大大提升了整体性能，具体流程如下所示：7基于寄存器的编程模型是指将数据从LocalTensor通过Reg搬运指令（Reg数据搬入）加载到寄存器中，进行复杂的数学计算Compute后通过Reg搬运指令（Reg数据搬出）搬出到LocalTensor中，所有的计算逻辑均在寄存器中完成，从而减少LocalTensor间的数据搬运，大大提升了整体性能，具体流程如下所示：
6 8 
7**图1** Reg矢量计算 9**图1** Reg矢量计算
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/register_data_types/MaskReg.md+2-0
@@ -36,6 +36,7 @@ MaskReg寄存器用于指示在计算过程中哪些元素参与计算，宽度
36![](../../../../figures/maskreg_process.png "MaskReg计算过程")36![](../../../../figures/maskreg_process.png "MaskReg计算过程")
37 37 
38## 函数原型<a name="section620mcpsimp"></a>38## 函数原型<a name="section620mcpsimp"></a>
39+ 
39- CreateMask接口40- CreateMask接口
40 ```cpp41 ```cpp
41 template <typename T, MaskPattern mode = MaskPattern::ALL, const RegTrait& regTrait = RegTraitNumOne>42 template <typename T, MaskPattern mode = MaskPattern::ALL, const RegTrait& regTrait = RegTraitNumOne>
@@ -267,6 +268,7 @@ __simd_callee__ inline void Add(U& dstReg, U& srcReg0, U& srcReg1, MaskReg& mask
267 ```268 ```
268 269 
269<!-- npu="910b,A3,950" id8 -->270<!-- npu="910b,A3,950" id8 -->
271+ 
270### NPU架构版本2201与NPU架构版本3510之间Mask机制对比272### NPU架构版本2201与NPU架构版本3510之间Mask机制对比
271 273 
272本节对比[NPU架构版本2201](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)在Normal模式的逐bit计算中Mask机制，与[NPU架构版本3510](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)基于Reg矢量计算中Mask机制的差异。274本节对比[NPU架构版本2201](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)在Normal模式的逐bit计算中Mask机制，与[NPU架构版本3510](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)基于Reg矢量计算中Mask机制的差异。
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/type_conversion/Cast.md+1-0
@@ -289,6 +289,7 @@ result[i][7 :0] = saturation(tmp[7 : 0]) according to control bit;
289```289```
290 290 
291### half转hifloat8_t转换规则291### half转hifloat8_t转换规则
292+ 
292```293```
293tmp2[15 : 0] = f16_src_data[i][15 : 0];294tmp2[15 : 0] = f16_src_data[i][15 : 0];
294Ev = tmp2[14 : 10] - 5'b01111, thr = {tmp2[0], 1’b1};295Ev = tmp2[14 : 10] - 5'b01111, thr = {tmp2[0], 1’b1};
Mdocs/zh/api/SIMD-API/basic_api/scalar_compute/Nop.md+1-0
@@ -26,6 +26,7 @@
26<!-- @ref: asc-devkit/res/docs/zh/api/SIMD-API/basic_api/scalar_compute/Nop_res.md#id1 -->26<!-- @ref: asc-devkit/res/docs/zh/api/SIMD-API/basic_api/scalar_compute/Nop_res.md#id1 -->
27 27 
28## 功能说明28## 功能说明
29+ 
29头文件路径为：`"basic_api/kernel_operator_utils_intf.h"`。30头文件路径为：`"basic_api/kernel_operator_utils_intf.h"`。
30 31 
31用户输入延迟的个数，实现延迟若干个cycle。32用户输入延迟的个数，实现延迟若干个cycle。
Mdocs/zh/api/SIMD-API/basic_api/sync_control/inter_core_sync/key_features.md+1-0
@@ -80,6 +80,7 @@
80- 在**3510架构**下占用模式4的flagId为1、17（AIC：1、17；AIV：1）。参见[DataCopy（GMToL1连续数据搬运）](../../cube_compute_ISASI/cube_compute_load/DataCopy_GMToL1_continuous.md)。80- 在**3510架构**下占用模式4的flagId为1、17（AIC：1、17；AIV：1）。参见[DataCopy（GMToL1连续数据搬运）](../../cube_compute_ISASI/cube_compute_load/DataCopy_GMToL1_continuous.md)。
81 81 
82## 各个核间同步控制实现原理<a id="sync_control_mode"></a>82## 各个核间同步控制实现原理<a id="sync_control_mode"></a>
83+ 
83如图1所示，本章节将配合时序图介绍CrossCoreSetFlag和CrossCoreWaitFlag配合使用时支持的四种同步控制模式各自实现的原理。84如图1所示，本章节将配合时序图介绍CrossCoreSetFlag和CrossCoreWaitFlag配合使用时支持的四种同步控制模式各自实现的原理。
84 85 
85- 模式0：AI Core核间的同步控制。对于AIC全核场景，同步所有的AIC核，直到所有的AIC核都执行到CrossCoreSetFlag时，CrossCoreWaitFlag后续的全部流水或者由模板参数pipe指定的流水（与NPU架构有关）中的指令才会执行；对于AIV全核场景，同步所有的AIV核，直到所有的AIV核都执行到CrossCoreSetFlag时，CrossCoreWaitFlag后续的全部流水或者由模板参数pipe指定的流水（与NPU架构有关）中的指令才会执行。86- 模式0：AI Core核间的同步控制。对于AIC全核场景，同步所有的AIC核，直到所有的AIC核都执行到CrossCoreSetFlag时，CrossCoreWaitFlag后续的全部流水或者由模板参数pipe指定的流水（与NPU架构有关）中的指令才会执行；对于AIV全核场景，同步所有的AIV核，直到所有的AIV核都执行到CrossCoreSetFlag时，CrossCoreWaitFlag后续的全部流水或者由模板参数pipe指定的流水（与NPU架构有关）中的指令才会执行。
Mdocs/zh/api/SIMD-API/c_api/cube_compute/asc_enable_hif8.md+1-0
@@ -55,6 +55,7 @@ PIPE_S
55- 本接口仅对矩阵乘加输入数据类型为`fp8_e4m3fn_t`×`fp8_e4m3fn_t`的场景生效，其他FP8数据类型组合（`fp8_e4m3fn_t`×`fp8_e5m2_t`、`fp8_e5m2_t`×`fp8_e4m3fn_t`、`fp8_e5m2_t`×`fp8_e5m2_t`）不支持FP8与HiF8模式选择，调用本接口不产生实际作用。55- 本接口仅对矩阵乘加输入数据类型为`fp8_e4m3fn_t`×`fp8_e4m3fn_t`的场景生效，其他FP8数据类型组合（`fp8_e4m3fn_t`×`fp8_e5m2_t`、`fp8_e5m2_t`×`fp8_e4m3fn_t`、`fp8_e5m2_t`×`fp8_e5m2_t`）不支持FP8与HiF8模式选择，调用本接口不产生实际作用。
56 56 
57<!-- npu="950" id8 -->57<!-- npu="950" id8 -->
58+ 
58## 调用示例59## 调用示例
59 60 
60将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。61将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
Mdocs/zh/api/SIMD-API/c_api/cube_compute/asc_mmad.md+2-0
@@ -67,6 +67,7 @@ $$
67## 函数原型67## 函数原型
68 68 
69<!-- npu="950" id12 -->69<!-- npu="950" id12 -->
70+ 
70### 函数原型（[NPU架构版本3510](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)）71### 函数原型（[NPU架构版本3510](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)）
71 72 
72**占位符形式：**73**占位符形式：**
@@ -143,6 +144,7 @@ __aicore__ inline void asc_mmad(__cc__ float* c_matrix,
143<!-- end id12 -->144<!-- end id12 -->
144 145 
145<!-- npu="A3,910b" id13 -->146<!-- npu="A3,910b" id13 -->
147+ 
146### 函数原型（[NPU架构版本2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)）148### 函数原型（[NPU架构版本2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)）
147 149 
148**占位符形式：**150**占位符形式：**
Mdocs/zh/api/SIMD-API/c_api/cube_compute/cube_compute_fractal_intro/aux_cube_fractal_format_details.md+2-0
@@ -1,6 +1,7 @@
1# 辅助矩阵分形格式详解1# 辅助矩阵分形格式详解
2 2 
3<!-- npu="950" id1 -->3<!-- npu="950" id1 -->
4+ 
4## MX系数矩阵5## MX系数矩阵
5 6 
6针对Ascend 950PR&950DT系列产品产品：7针对Ascend 950PR&950DT系列产品产品：
@@ -71,6 +72,7 @@ MX矩阵乘法为带有量化系数X的矩阵乘法，即左矩阵和右矩阵
71<!-- end id1 -->72<!-- end id1 -->
72 73 
73<!-- npu="A3,910b" id2 -->74<!-- npu="A3,910b" id2 -->
75+ 
74## 4选2稀疏索引矩阵76## 4选2稀疏索引矩阵
75 77 
76针对Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品：78针对Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品：
Mdocs/zh/api/SIMD-API/c_api/cube_compute/cube_compute_fractal_intro/key_fractal_format_details.md+1-0
@@ -61,6 +61,7 @@
61<!-- end id5 -->61<!-- end id5 -->
62 62 
63<!-- npu="910b,A3" id6 -->63<!-- npu="910b,A3" id6 -->
64+ 
64## Zz格式（FRACTAL_Zz）65## Zz格式（FRACTAL_Zz）
65 66 
66- **定义：** 大Z外部row major（行主序）+ 小z内部row major（行主序），见下图。67- **定义：** 大Z外部row major（行主序）+ 小z内部row major（行主序），见下图。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_copy_gm2l0a.md+1-0
@@ -72,6 +72,7 @@ PIPE_MTE2
72- 操作数地址重叠约束请参考[通用地址重叠约束](../general_description_and_constraints.md#通用地址重叠约束)。72- 操作数地址重叠约束请参考[通用地址重叠约束](../general_description_and_constraints.md#通用地址重叠约束)。
73 73 
74## 调用示例74## 调用示例
75+ 
75```cpp76```cpp
76__ca__ half dst[256];77__ca__ half dst[256];
77// src表示源操作数的起始地址78// src表示源操作数的起始地址
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_set_l0c_copy_config.md+1-0
@@ -85,6 +85,7 @@ PIPE_S
85- 本接口的配置会持续生效，直至再次调用本接口覆盖。不同搬出任务使用的配置不同时，需在对应搬出接口执行前重新配置，每次调用均会同时覆盖三个参数。85- 本接口的配置会持续生效，直至再次调用本接口覆盖。不同搬出任务使用的配置不同时，需在对应搬出接口执行前重新配置，每次调用均会同时覆盖三个参数。
86 86 
87<!-- npu="950" id12 -->87<!-- npu="950" id12 -->
88+ 
88## 调用示例89## 调用示例
89 90 
90将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。91将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_set_l0c_copy_lrelu_alpha.md+1-0
@@ -62,6 +62,7 @@ PIPE_S
62- 调用后配置会持续生效，直至再次调用本接口覆盖，不同搬出任务使用的系数不同时，需在对应搬出接口执行前重新配置。62- 调用后配置会持续生效，直至再次调用本接口覆盖，不同搬出任务使用的系数不同时，需在对应搬出接口执行前重新配置。
63 63 
64<!-- npu="950" id8 -->64<!-- npu="950" id8 -->
65+ 
65## 调用示例66## 调用示例
66 67 
67将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。68将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_set_l0c_copy_relu_alpha.md+1-0
@@ -62,6 +62,7 @@ PIPE_S
62- 调用后配置会持续生效，直至再次调用本接口覆盖，不同搬出任务使用的系数不同时，需在对应搬出接口执行前重新配置。62- 调用后配置会持续生效，直至再次调用本接口覆盖，不同搬出任务使用的系数不同时，需在对应搬出接口执行前重新配置。
63 63 
64<!-- npu="950" id8 -->64<!-- npu="950" id8 -->
65+ 
65## 调用示例66## 调用示例
66 67 
67将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。68将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_set_l13d_fmatrix.md+1-0
@@ -75,6 +75,7 @@ PIPE_S
75- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。75- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。
76 76 
77<!-- npu="950" id8 -->77<!-- npu="950" id8 -->
78+ 
78## 调用示例79## 调用示例
79 80 
80将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[`__NPU_ARCH__`](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。81将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[`__NPU_ARCH__`](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_set_l13d_fmatrix_b.md+1-0
@@ -76,6 +76,7 @@ PIPE_S
76- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。76- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。
77 77 
78<!-- npu="950" id8 -->78<!-- npu="950" id8 -->
79+ 
79## 调用示例80## 调用示例
80 81 
81将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[`__NPU_ARCH__`](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。82将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[`__NPU_ARCH__`](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_set_l13d_rpt.md+1-0
@@ -74,6 +74,7 @@ PIPE_S
74- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。74- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。
75 75 
76<!-- npu="950" id8 -->76<!-- npu="950" id8 -->
77+ 
77## 调用示例78## 调用示例
78 79 
79将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[`__NPU_ARCH__`](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。80将代码保存为`example.asc`后，可通过`bisheng`命令编译运行，其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构，具体产品与NPU架构的映射关系请参考[`__NPU_ARCH__`](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_set_l13d_rpt_b.md+1-0
@@ -69,6 +69,7 @@ PIPE_S
69- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。69- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。
70 70 
71<!-- npu="950" id8 -->71<!-- npu="950" id8 -->
72+ 
72## 调用示例73## 调用示例
73 74 
74以下示例使用`asc_copy_l12l0b`将2×2、16通道的全零Feature Map扩展为带一圈padding的4×4矩阵，并自动转置写入L0B Buffer。B侧padding值设置为2，`f_matrix_ctrl`设置为`true`。使用全1的L0A矩阵进行MMAD后，每个输出累加4×4区域中的12个padding位置，期望值为`12 × 2 = 24`。75以下示例使用`asc_copy_l12l0b`将2×2、16通道的全零Feature Map扩展为带一圈padding的4×4矩阵，并自动转置写入L0B Buffer。B侧padding值设置为2，`f_matrix_ctrl`设置为`true`。使用全1的L0A矩阵进行MMAD后，每个输出累加4×4区域中的12个padding位置，期望值为`12 × 2 = 24`。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_set_l3d_rpt_b.md+1-0
@@ -72,6 +72,7 @@ PIPE_S
72- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。72- 调用后配置会持续生效，直至再次调用本接口覆盖，不同3D搬运任务使用的配置不同时，需在对应搬运接口执行前重新配置。
73 73 
74<!-- npu="950" id8 -->74<!-- npu="950" id8 -->
75+ 
75## 调用示例76## 调用示例
76 77 
77`asc_set_l3d_rpt_b`需要调用者自行拼装64位配置。相同功能推荐使用参数化接口[asc_set_l13d_rpt_b](asc_set_l13d_rpt_b.md)，以下示例使用该接口配置右矩阵的repeat参数。78`asc_set_l3d_rpt_b`需要调用者自行拼装64位配置。相同功能推荐使用参数化接口[asc_set_l13d_rpt_b](asc_set_l13d_rpt_b.md)，以下示例使用该接口配置右矩阵的repeat参数。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_input_move_constraint.md+2-0
@@ -34,6 +34,7 @@ asc_copy_l12l0b_sparse主要用于搬运以512Byte为单位存放的压缩权重
34矩阵计算过程中，常用分形支持总结如下表：34矩阵计算过程中，常用分形支持总结如下表：
35 35 
36<!-- npu="910b,A3" id7 -->36<!-- npu="910b,A3" id7 -->
37+ 
37### 针对Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品38### 针对Atlas A2训练系列产品/Atlas A2推理系列产品和Atlas A3训练系列产品/Atlas A3推理系列产品
38 39 
39常用分形支持情况如下，供开发者参考：40常用分形支持情况如下，供开发者参考：
@@ -119,6 +120,7 @@ asc_copy_l12l0b_sparse主要用于搬运以512Byte为单位存放的压缩权重
119<!-- end id7 -->120<!-- end id7 -->
120 121 
121<!-- npu="950" id8 -->122<!-- npu="950" id8 -->
123+ 
122### 针对Ascend 950PR&950DT系列产品124### 针对Ascend 950PR&950DT系列产品
123 125 
124数据搬运（GM-\>L1 Buffer）支持使能随路进行DN到Nz转换，一定会改变分形排布。126数据搬运（GM-\>L1 Buffer）支持使能随路进行DN到Nz转换，一定会改变分形排布。
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_compute_load_overall_description.md+1-0
@@ -48,6 +48,7 @@
48<!-- end id2 -->48<!-- end id2 -->
49 49 
50<!-- npu="950" id5 -->50<!-- npu="950" id5 -->
51+ 
51## UB->L1 Buffer的接口分类与使用场景52## UB->L1 Buffer的接口分类与使用场景
52 53 
53**表4** UB->L1 Buffer接口54**表4** UB->L1 Buffer接口
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_clear_ar_spr_deprecated.md+1-0
@@ -1,4 +1,5 @@
1# asc_clear_ar_spr（废弃）1# asc_clear_ar_spr（废弃）
2+ 
2## 产品支持情况3## 产品支持情况
3 4 
4<!-- npu="950" id1 -->5<!-- npu="950" id1 -->
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_copy_gm2l0a_sync_deprecated.md+1-0
@@ -74,6 +74,7 @@ PIPE_MTE2
74- 操作数地址重叠约束请参考[通用地址重叠约束](../general_description_and_constraints.md#通用地址重叠约束)。74- 操作数地址重叠约束请参考[通用地址重叠约束](../general_description_and_constraints.md#通用地址重叠约束)。
75 75 
76## 调用示例76## 调用示例
77+ 
77```cpp78```cpp
78__ca__ half dst[256];79__ca__ half dst[256];
79// src表示源操作数的起始地址80// src表示源操作数的起始地址
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_gather_datablock_sync_deprecated.md+1-0
@@ -63,6 +63,7 @@
63PIPE_V63PIPE_V
64 64 
65## 约束说明65## 约束说明
66+ 
66- src_offset取值应保证src元素类型位宽对齐。67- src_offset取值应保证src元素类型位宽对齐。
67- src加上src_offset偏移后的地址不能超出Unified Buffer（UB）大小数据的范围。68- src加上src_offset偏移后的地址不能超出Unified Buffer（UB）大小数据的范围。
68- dst、src的起始地址需要32字节对齐。69- dst、src的起始地址需要32字节对齐。
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_int322uint8_deprecated.md+1-0
@@ -1,4 +1,5 @@
1# asc_int322uint8（废弃）1# asc_int322uint8（废弃）
2+ 
2## 产品支持情况3## 产品支持情况
3 4 
4<!-- npu="950" id1 -->5<!-- npu="950" id1 -->
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_max_sync_deprecated.md+1-0
@@ -48,6 +48,7 @@ __aicore__ inline void asc_max_sync(__ubuf__ float* dst, __ubuf__ float* src0, _
48```48```
49 49 
50## 参数说明50## 参数说明
51+ 
51**表1** 参数说明52**表1** 参数说明
52 53 
53| 参数名 | 输入/输出 | 描述 |54| 参数名 | 输入/输出 | 描述 |
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_mmad_deprecated.md+2-0
@@ -82,6 +82,7 @@ K0的取值为`32B / sizeof(dtype)`，`dtype`为矩阵的数据类型。当数
82## 函数原型82## 函数原型
83 83 
84<!-- npu="950" id12 -->84<!-- npu="950" id12 -->
85+ 
85### 函数原型（[NPU架构版本3510](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)）86### 函数原型（[NPU架构版本3510](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)）
86 87 
87**占位符形式：**88**占位符形式：**
@@ -145,6 +146,7 @@ __aicore__ inline void asc_mmad(__cc__ float* c_matrix,
145<!-- end id12 -->146<!-- end id12 -->
146 147 
147<!-- npu="A3,910b" id20 -->148<!-- npu="A3,910b" id20 -->
149+ 
148### 函数原型（[NPU架构版本2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)）150### 函数原型（[NPU架构版本2201](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md)）
149 151 
150- `asc_mmad`及`asc_mmad_sync`：152- `asc_mmad`及`asc_mmad_sync`：
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_set_l0c2gm_config_deprecated.md+2-0
@@ -1,4 +1,5 @@
1# asc_set_l0c2gm_config（废弃）1# asc_set_l0c2gm_config（废弃）
2+ 
2## 产品支持情况3## 产品支持情况
3 4 
4<!-- npu="950" id1 -->5<!-- npu="950" id1 -->
@@ -46,6 +47,7 @@ __aicore__ inline void asc_set_l0c2gm_config(uint64_t relu_pre, uint64_t quant_p
46| relu_pre | 输入 | ReLU操作前矢量的起始地址。|47| relu_pre | 输入 | ReLU操作前矢量的起始地址。|
47| quant_pre | 输入 | 量化操作前矢量的起始地址。|48| quant_pre | 输入 | 量化操作前矢量的起始地址。|
48| enable_unit_flag | 输入 | 是否启用unit_flag。unit_flag是一种矩阵计算指令和矩阵搬运指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出，该功能不适用于L0C Buffer累加的场景。|49| enable_unit_flag | 输入 | 是否启用unit_flag。unit_flag是一种矩阵计算指令和矩阵搬运指令细粒度的并行，开启该功能后，硬件每计算完一个分形，计算结果就会被搬出，该功能不适用于L0C Buffer累加的场景。|
50+ 
49## 返回值说明51## 返回值说明
50 52 
51无53无
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_set_l0c2gm_quant_post_deprecated.md+1-0
@@ -72,6 +72,7 @@ PIPE_S
72量化参数不能为inf/nan和非规格化数。72量化参数不能为inf/nan和非规格化数。
73 73 
74## 调用示例74## 调用示例
75+ 
75```cpp76```cpp
76const uint64_t MASK_8_0 = 0x1FF;77const uint64_t MASK_8_0 = 0x1FF;
77const uint64_t MASK_BIT9 = 0x200;78const uint64_t MASK_BIT9 = 0x200;
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_set_l12l0_padding_val_deprecated.md+1-0
@@ -67,6 +67,7 @@ PIPE_S
67无67无
68 68 
69## 调用示例69## 调用示例
70+ 
70```cpp71```cpp
71uint64_t config = 0;72uint64_t config = 0;
72asc_set_l12l0_padding_val(config);73asc_set_l12l0_padding_val(config);
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_sqrt_sync_deprecated.md+1-0
@@ -62,6 +62,7 @@ $$
62PIPE_V62PIPE_V
63 63 
64## 约束说明64## 约束说明
65+ 
65- src只支持设置为正数。66- src只支持设置为正数。
66- 操作数地址重叠约束请参考[通用地址重叠约束](../general_description_and_constraints.md#通用地址重叠约束)。67- 操作数地址重叠约束请参考[通用地址重叠约束](../general_description_and_constraints.md#通用地址重叠约束)。
67- dst、src的起始地址需要32字节对齐。68- dst、src的起始地址需要32字节对齐。
Mdocs/zh/api/SIMD-API/c_api/reg_compute/key_features.md+3-0
@@ -5,6 +5,7 @@
5Reg矢量计算架构中，Vector Function（VF）是实现高性能向量计算的核心载体。VF函数中可以包含最多四层嵌套循环，每层循环中还可以包含多个串行循环。VF循环对控制结构的支持有限，仅支持for循环和条件判断，不支持switch、do-while和while-do等其他控制结构。当VF函数中的循环满足Hardware Loop编码规范会被编译器优化为Hardware Loop，提升整体的编码性能，否则它的循环逻辑会由迭代变量和条件判断语句构成Software Loop，无法开启VF循环优化。5Reg矢量计算架构中，Vector Function（VF）是实现高性能向量计算的核心载体。VF函数中可以包含最多四层嵌套循环，每层循环中还可以包含多个串行循环。VF循环对控制结构的支持有限，仅支持for循环和条件判断，不支持switch、do-while和while-do等其他控制结构。当VF函数中的循环满足Hardware Loop编码规范会被编译器优化为Hardware Loop，提升整体的编码性能，否则它的循环逻辑会由迭代变量和条件判断语句构成Software Loop，无法开启VF循环优化。
6 6 
7### Hardware Loop编码规范7### Hardware Loop编码规范
8+ 
8- Hardware Loop最多能支持的嵌套层数为4层；支持串行的Hardware Loop循环。9- Hardware Loop最多能支持的嵌套层数为4层；支持串行的Hardware Loop循环。
9- 迭代变量类型<br>VF内所有Loop的迭代变量必须是uint16_t类型。10- 迭代变量类型<br>VF内所有Loop的迭代变量必须是uint16_t类型。
10- 起始值与步长<br>循环起始值从0开始。<br>每次迭代的步长必须是递增1。11- 起始值与步长<br>循环起始值从0开始。<br>每次迭代的步长必须是递增1。
@@ -13,6 +14,7 @@ Reg矢量计算架构中，Vector Function（VF）是实现高性能向量计算
13- 若要利用外层循环的计数作为循环边界，将外层循环计数器赋值给标量后作为内层循环边界。14- 若要利用外层循环的计数作为循环边界，将外层循环计数器赋值给标量后作为内层循环边界。
14 15 
15## 指令双发优化16## 指令双发优化
17+ 
16Reg矢量计算接口指令单双发的情况请参考[Reg矢量计算接口指令单双发汇总](../../../appendix/reg_vector_compute_interface_instruction_single_dual_issue_summary.md)。18Reg矢量计算接口指令单双发的情况请参考[Reg矢量计算接口指令单双发汇总](../../../appendix/reg_vector_compute_interface_instruction_single_dual_issue_summary.md)。
17 19 
18指令双发优化的性能优化详细内容请参考[算子实践参考-指令双发优化](../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/dual_issue_optimization.md)。20指令双发优化的性能优化详细内容请参考[算子实践参考-指令双发优化](../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/dual_issue_optimization.md)。
@@ -32,6 +34,7 @@ VF并不是写的越长，把所有运算都放在一个for循环内就好，需
32如果循环内存在依赖关系过多的指令，指令无法并发执行，即无法启用双发特性。可以通过展开循环，提升指令双发能力，贴近硬件乱序执行的特性；减少指令因为寄存器资源未到位而产生的等待。34如果循环内存在依赖关系过多的指令，指令无法并发执行，即无法启用双发特性。可以通过展开循环，提升指令双发能力，贴近硬件乱序执行的特性；减少指令因为寄存器资源未到位而产生的等待。
33 35 
34## VF融合36## VF融合
37+ 
35VF融合是将代码中多个VF函数融合成一个VF函数，有效提升性能。VF融合特性是编译器优化特性，通过编译选项--cce-simd-vf-fusion=true开启，VF自动融合会借助Loop Fuse算法，将VF转换成Loop形态，然后将控制流等价（Control-Flow-Equivalent）的VF进行融合，最后将VF进行还原。编译器首先会做融合前的合法性检查，判断两个VF是否等价，Main侧中间代码是否能在VF内执行以及融合后是否可产生正收益。如果满足VF融合条件，编译器会自动执行VF融合优化，为保证融合后的VF执行逻辑与语义与融合前一致，会在原来两个VF之间保守地插入同步指令，编译器还会尝试外提、合并融合后的VF中的指令，对VF代码进行优化。融合策略是能融尽融，用户按照符合融合的合法性检查的模式进行编码，可以增加VF融合的机会。用户也可以参考融合原理手动进行融合优化。38VF融合是将代码中多个VF函数融合成一个VF函数，有效提升性能。VF融合特性是编译器优化特性，通过编译选项--cce-simd-vf-fusion=true开启，VF自动融合会借助Loop Fuse算法，将VF转换成Loop形态，然后将控制流等价（Control-Flow-Equivalent）的VF进行融合，最后将VF进行还原。编译器首先会做融合前的合法性检查，判断两个VF是否等价，Main侧中间代码是否能在VF内执行以及融合后是否可产生正收益。如果满足VF融合条件，编译器会自动执行VF融合优化，为保证融合后的VF执行逻辑与语义与融合前一致，会在原来两个VF之间保守地插入同步指令，编译器还会尝试外提、合并融合后的VF中的指令，对VF代码进行优化。融合策略是能融尽融，用户按照符合融合的合法性检查的模式进行编码，可以增加VF融合的机会。用户也可以参考融合原理手动进行融合优化。
36 39 
37VF融合的性能优化详细内容请参考[算子实践参考-VF融合优化](../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/vf_fusion_optimization.md)。40VF融合的性能优化详细内容请参考[算子实践参考-VF融合优化](../../../../guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/vf_fusion_optimization.md)。
Mdocs/zh/api/SIMD-API/c_api/reg_compute/overview.md+2-0
@@ -1,7 +1,9 @@
1# 概述1# 概述
2+ 
2Reg矢量计算C API提供面向Reg矢量计算架构的寄存器级矢量计算能力，用户可以通过该API直接对芯片中涉及Vector计算的寄存器进行操作，实现更大的灵活性和更好的性能。Reg矢量计算C API与Memory矢量计算C API功能相似，但与Memory矢量计算C API的输入和输出数据通常位于Unified Buffer（UB）不同，Reg矢量计算C API的输入或输出数据可使用矢量数据寄存器`vector_<dtype>`（`dtype`为占位符，表示`half`、`float`等C API支持的数据类型）和掩码寄存器`vector_bool`。对于计算类API，其功能是从给定的寄存器获取数据，进行计算，并将结果保存在给定的寄存器。对于搬运类API，其功能是实现UB和寄存器的数据搬运。由此可见，Reg矢量计算C API相较于Memory矢量计算C API，将数据搬运和Reg计算过程交给用户自主控制，从而实现更大的开发自由度。3Reg矢量计算C API提供面向Reg矢量计算架构的寄存器级矢量计算能力，用户可以通过该API直接对芯片中涉及Vector计算的寄存器进行操作，实现更大的灵活性和更好的性能。Reg矢量计算C API与Memory矢量计算C API功能相似，但与Memory矢量计算C API的输入和输出数据通常位于Unified Buffer（UB）不同，Reg矢量计算C API的输入或输出数据可使用矢量数据寄存器`vector_<dtype>`（`dtype`为占位符，表示`half`、`float`等C API支持的数据类型）和掩码寄存器`vector_bool`。对于计算类API，其功能是从给定的寄存器获取数据，进行计算，并将结果保存在给定的寄存器。对于搬运类API，其功能是实现UB和寄存器的数据搬运。由此可见，Reg矢量计算C API相较于Memory矢量计算C API，将数据搬运和Reg计算过程交给用户自主控制，从而实现更大的开发自由度。
3 4 
4## 如何使用Reg矢量计算API5## 如何使用Reg矢量计算API
6+ 
5基于寄存器的编程模型是指将数据从UB通过Reg搬运指令（Reg数据搬入）加载到寄存器中，进行复杂的数学计算后通过Reg搬运指令（Reg数据搬出）搬出到UB中，所有的计算逻辑均在寄存器中完成，从而减少UB中间数据的反复读写，大大提升整体性能，具体流程如下所示：7基于寄存器的编程模型是指将数据从UB通过Reg搬运指令（Reg数据搬入）加载到寄存器中，进行复杂的数学计算后通过Reg搬运指令（Reg数据搬出）搬出到UB中，所有的计算逻辑均在寄存器中完成，从而减少UB中间数据的反复读写，大大提升整体性能，具体流程如下所示：
6 8 
7**图1** Reg矢量计算9**图1** Reg矢量计算
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_arith/asc_sub.md+1-0
@@ -110,6 +110,7 @@ __simd_callee__ inline void asc_sub(vector_bool& carry,
110 vector_int32_t src1,110 vector_int32_t src1,
111 vector_bool mask)111 vector_bool mask)
112```112```
113+ 
113## 参数说明114## 参数说明
114 115 
115### 无借位输出减法116### 无借位输出减法
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/rounding_mode.md+12-0
@@ -26,6 +26,7 @@ Ascend C API支持以下舍入模式：
26### 舍入模式详细说明26### 舍入模式详细说明
27 27 
28#### RINT（四舍六入五成双）28#### RINT（四舍六入五成双）
29+ 
29也称为"银行家舍入法"，是一种统计学上最无偏的舍入方法。30也称为"银行家舍入法"，是一种统计学上最无偏的舍入方法。
30- 当小数部分小于0.5时，向下舍入31- 当小数部分小于0.5时，向下舍入
31- 当小数部分大于0.5时，向上舍入32- 当小数部分大于0.5时，向上舍入
@@ -41,6 +42,7 @@ Ascend C API支持以下舍入模式：
41```42```
42 43 
43#### ROUND（四舍五入）44#### ROUND（四舍五入）
45+ 
44最常用的舍入方法，符合日常习惯。46最常用的舍入方法，符合日常习惯。
45- 当小数部分小于0.5时，向下舍入47- 当小数部分小于0.5时，向下舍入
46- 当小数部分大于等于0.5时，向上舍入48- 当小数部分大于等于0.5时，向上舍入
@@ -54,6 +56,7 @@ Ascend C API支持以下舍入模式：
54```56```
55 57 
56#### FLOOR（向下取整）58#### FLOOR（向下取整）
59+ 
57始终向负无穷方向舍入，即使结果变小。60始终向负无穷方向舍入，即使结果变小。
58 61 
59**示例：**62**示例：**
@@ -63,6 +66,7 @@ Ascend C API支持以下舍入模式：
63```66```
64 67 
65#### CEIL（向上取整）68#### CEIL（向上取整）
69+ 
66始终向正无穷方向舍入，即使结果变大。70始终向正无穷方向舍入，即使结果变大。
67 71 
68**示例：**72**示例：**
@@ -72,6 +76,7 @@ Ascend C API支持以下舍入模式：
72```76```
73 77 
74#### TRUNC（向零截断）78#### TRUNC（向零截断）
79+ 
75直接截断小数部分，向零方向舍入。80直接截断小数部分，向零方向舍入。
76 81 
77**示例：**82**示例：**
@@ -81,6 +86,7 @@ Ascend C API支持以下舍入模式：
81```86```
82 87 
83#### ODD（奇数舍入）88#### ODD（奇数舍入）
89+ 
84舍入到最近的奇数，主要用于某些特殊的数值计算场景。90舍入到最近的奇数，主要用于某些特殊的数值计算场景。
85 91 
86**示例：**92**示例：**
@@ -103,6 +109,7 @@ Ascend C API支持以下舍入模式：
103在非饱和模式下，当数据超出目标类型范围时：109在非饱和模式下，当数据超出目标类型范围时：
104 110 
105#### 整型转换111#### 整型转换
112+ 
106- **结果被截断**：保留最低有效位（LSB），高位直接丢弃113- **结果被截断**：保留最低有效位（LSB），高位直接丢弃
107- **二进制截断示例**：将int32转换为int8时，只保留低8位114- **二进制截断示例**：将int32转换为int8时，只保留低8位
108 115 
@@ -125,6 +132,7 @@ Ascend C API支持以下舍入模式：
125```132```
126 133 
127#### 浮点数转换134#### 浮点数转换
135+ 
128- **超出范围**：返回输出类型的对应符号inf值（正/负无穷）136- **超出范围**：返回输出类型的对应符号inf值（正/负无穷）
129- **±inf输入**：返回±inf137- **±inf输入**：返回±inf
130- **NaN输入**：返回NaN138- **NaN输入**：返回NaN
@@ -142,6 +150,7 @@ Ascend C API支持以下舍入模式：
142```150```
143 151 
144#### 浮点转整型152#### 浮点转整型
153+ 
145- **超出范围**：结果被截断至目标格式的数据宽度154- **超出范围**：结果被截断至目标格式的数据宽度
146- **±inf输入**：返回输出类型的最值155- **±inf输入**：返回输出类型的最值
147- **NaN输入**：返回0156- **NaN输入**：返回0
@@ -163,6 +172,7 @@ Ascend C API支持以下舍入模式：
163在饱和模式下，当数据超出目标类型范围时，结果被限制在目标类型的可表示范围内：172在饱和模式下，当数据超出目标类型范围时，结果被限制在目标类型的可表示范围内：
164 173 
165#### 整型转换174#### 整型转换
175+ 
166- **超出范围**：返回目标类型的最大值或最小值176- **超出范围**：返回目标类型的最大值或最小值
167- **保持数值安全**：不会出现因截断导致的意外结果177- **保持数值安全**：不会出现因截断导致的意外结果
168 178 
@@ -191,6 +201,7 @@ Ascend C API支持以下舍入模式：
191```201```
192 202 
193#### 浮点数转换203#### 浮点数转换
204+ 
194- **超出范围**：返回输出类型的对应最值205- **超出范围**：返回输出类型的对应最值
195- **±inf输入**：返回输出类型的对应最值206- **±inf输入**：返回输出类型的对应最值
196- **NaN输入**：返回0207- **NaN输入**：返回0
@@ -211,6 +222,7 @@ Ascend C API支持以下舍入模式：
211```222```
212 223 
213#### 浮点转整型224#### 浮点转整型
225+ 
214- **超出范围**：返回输出类型的对应最值226- **超出范围**：返回输出类型的对应最值
215- **±inf输入**：返回输出类型的对应最值227- **±inf输入**：返回输出类型的对应最值
216- **NaN输入**：返回0228- **NaN输入**：返回0
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_copy/asc_copy.md+1-0
@@ -56,6 +56,7 @@ __simd_callee__ inline void asc_copy(vector_<dtype>& dst,
56```56```
57 57 
58#### dtype支持数据类型58#### dtype支持数据类型
59+ 
59`dtype`取值为：`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。60`dtype`取值为：`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。
60 61 
61#### 函数原型典型示例62#### 函数原型典型示例
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_copy/asc_mask.md+1-0
@@ -49,6 +49,7 @@ __simd_callee__ inline vector_<dtype> asc_mask(vector_<dtype> src,
49```49```
50 50 
51#### dtype支持数据类型51#### dtype支持数据类型
52+ 
52`dtype`取值为：`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。53`dtype`取值为：`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`half`、`bfloat16_t`、`int32_t`、`uint32_t`、`float`。
53 54 
54#### 函数原型典型示例55#### 函数原型典型示例
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_gather/asc_gather.md+1-0
@@ -52,6 +52,7 @@ __simd_callee__ inline vector_<dtype> asc_gather(vector_<dtype> src,
52```52```
53 53 
54### 支持数据类型列表54### 支持数据类型列表
55+ 
55**表1** 寄存器源收集模式支持数据类型列表56**表1** 寄存器源收集模式支持数据类型列表
56 57 
57| dtype | index_dtype |58| dtype | index_dtype |
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_logic/asc_xor.md+1-0
@@ -63,6 +63,7 @@ __simd_callee__ inline void asc_xor(vector_<dtype>& dst,
63```63```
64 64 
65#### dtype支持数据类型65#### dtype支持数据类型
66+ 
66`dtype`支持的数据类型：`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`int32_t`、`uint32_t`。67`dtype`支持的数据类型：`int8_t`、`uint8_t`、`int16_t`、`uint16_t`、`int32_t`、`uint32_t`。
67 68 
68#### 函数原型典型示例69#### 函数原型典型示例
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_permute_sel/asc_select.md+1-0
@@ -23,6 +23,7 @@
23<!-- npu="910" id7 -->23<!-- npu="910" id7 -->
24- Atlas训练系列产品：不支持24- Atlas训练系列产品：不支持
25<!-- end id7 -->25<!-- end id7 -->
26+ 
26## 功能说明27## 功能说明
27 28 
28头文件路径为：`"c_api/reg_compute/compute/reg_permute_sel.h"`。29头文件路径为：`"c_api/reg_compute/compute/reg_permute_sel.h"`。
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_permute_sel/asc_unpack.md+1-0
@@ -72,6 +72,7 @@ __simd_callee__ inline void asc_unpack_upper(vector_<dtype_dst>& dst,
72```72```
73 73 
74#### dtype支持数据类型74#### dtype支持数据类型
75+ 
75dtype_src与dtype_dst支持的数据类型对如下：76dtype_src与dtype_dst支持的数据类型对如下：
76 77 
77| dtype_src | dtype_dst |78| dtype_src | dtype_dst |
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_permute_sel/asc_unsqueeze.md+1-0
@@ -55,6 +55,7 @@ __simd_callee__ inline void asc_unsqueeze(vector_<dtype>& dst,
55```55```
56 56 
57#### dtype支持数据类型57#### dtype支持数据类型
58+ 
58dtype支持的数据类型：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t。返回值类型接口通过函数名后缀区分数据类型，对应关系如下：59dtype支持的数据类型：int8_t、uint8_t、int16_t、uint16_t、int32_t、uint32_t。返回值类型接口通过函数名后缀区分数据类型，对应关系如下：
59 60 
60| 函数名后缀 | 数据类型 |61| 函数名后缀 | 数据类型 |
Mdocs/zh/api/SIMD-API/c_api/reg_compute/ub_gather/asc_gather.md+1-0
@@ -54,6 +54,7 @@ __simd_callee__ inline vector_<dst_dtype> asc_gather(__ubuf__ <src_dtype>* src,
54```54```
55 55 
56### 数据类型组合列表56### 数据类型组合列表
57+ 
57**表1** UB源收集模式支持数据类型组合列表58**表1** UB源收集模式支持数据类型组合列表
58 59 
59| dst_dtype | src_dtype | index_dtype |60| dst_dtype | src_dtype | index_dtype |
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_arith/asc_max.md+1-0
@@ -63,6 +63,7 @@ __aicore__ inline void asc_max(__ubuf__ float* dst, __ubuf__ float* src0, __ubuf
63```63```
64 64 
65## 参数说明65## 参数说明
66+ 
66**表1** 参数说明67**表1** 参数说明
67 68 
68| 参数名 | 输入/输出 | 描述 |69| 参数名 | 输入/输出 | 描述 |
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_arith/asc_sqrt.md+1-0
@@ -71,6 +71,7 @@ $$
71PIPE_V71PIPE_V
72 72 
73## 约束说明73## 约束说明
74+ 
74- src只支持设置为正数。75- src只支持设置为正数。
75- 操作数地址重叠约束请参考[通用地址重叠约束](../../general_description_and_constraints.md#通用地址重叠约束)。76- 操作数地址重叠约束请参考[通用地址重叠约束](../../general_description_and_constraints.md#通用地址重叠约束)。
76- dst、src的起始地址需要32字节对齐。77- dst、src的起始地址需要32字节对齐。
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_compare/asc_ge.md+1-0
@@ -35,6 +35,7 @@ dst = src0 >= src1
35$$35$$
36 36 
37## 函数原型37## 函数原型
38+ 
38- 高维切分计算39- 高维切分计算
39 40 
40 ```cpp41 ```cpp
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_fused/asc_leakyrelu.md+1-0
@@ -39,6 +39,7 @@ src_i ,\quad src_i>0\\
39$$39$$
40 40 
41## 函数原型41## 函数原型
42+ 
42- 前n个数据计算43- 前n个数据计算
43 44 
44 ```cpp45 ```cpp
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_gather/asc_gather_datablock.md+1-0
@@ -61,6 +61,7 @@
61PIPE_V61PIPE_V
62 62 
63## 约束说明63## 约束说明
64+ 
64- src_offset取值应保证src元素类型位宽对齐。65- src_offset取值应保证src元素类型位宽对齐。
65- src加上src_offset偏移后的地址不能超出Unified Buffer（UB）大小数据的范围。66- src加上src_offset偏移后的地址不能超出Unified Buffer（UB）大小数据的范围。
66- dst、src的起始地址需要32字节对齐。67- dst、src的起始地址需要32字节对齐。
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_sort/asc_mrgsort4.md+1-0
@@ -120,6 +120,7 @@ PIPE_V
120<!-- end id12 -->120<!-- end id12 -->
121 121 
122<!-- npu="950" id8 -->122<!-- npu="950" id8 -->
123+ 
123## 调用示例124## 调用示例
124 125 
125以下示例将256个乱序float类型score及其index分别作为输入，通过[asc_bitsort](asc_bitsort.md)与`asc_mrgsort4`配合完成全局降序排序，每个输出数据单元包含一个score及其对应的index，占8字节，总体计算流程为：126以下示例将256个乱序float类型score及其index分别作为输入，通过[asc_bitsort](asc_bitsort.md)与`asc_mrgsort4`配合完成全局降序排序，每个输出数据单元包含一个score及其对应的index，占8字节，总体计算流程为：
Mdocs/zh/api/SIMD-API/tensor_api/reg_vector_compute/compare_and_select/operator_compare.md+6-0
@@ -42,6 +42,7 @@
42## 函数原型42## 函数原型
43 43 
44<a id="operator-eq"></a>44<a id="operator-eq"></a>
45+ 
45### operator==46### operator==
46 47 
47```cpp48```cpp
@@ -59,6 +60,7 @@ __simd_callee__ inline reg_tensor<bool> operator==(
59```60```
60 61 
61<a id="operator-ne"></a>62<a id="operator-ne"></a>
63+ 
62### operator!=64### operator!=
63 65 
64```cpp66```cpp
@@ -76,6 +78,7 @@ __simd_callee__ inline reg_tensor<bool> operator!=(
76```78```
77 79 
78<a id="operator-lt"></a>80<a id="operator-lt"></a>
81+ 
79### operator<82### operator<
80 83 
81```cpp84```cpp
@@ -93,6 +96,7 @@ __simd_callee__ inline reg_tensor<bool> operator<(
93```96```
94 97 
95<a id="operator-le"></a>98<a id="operator-le"></a>
99+ 
96### operator<=100### operator<=
97 101 
98```cpp102```cpp
@@ -110,6 +114,7 @@ __simd_callee__ inline reg_tensor<bool> operator<=(
110```114```
111 115 
112<a id="operator-gt"></a>116<a id="operator-gt"></a>
117+ 
113### operator>118### operator>
114 119 
115```cpp120```cpp
@@ -127,6 +132,7 @@ __simd_callee__ inline reg_tensor<bool> operator>(
127```132```
128 133 
129<a id="operator-ge"></a>134<a id="operator-ge"></a>
135+ 
130### operator>=136### operator>=
131 137 
132```cpp138```cpp
Mdocs/zh/api/Utils-API/tuning_interface/asc_dump.md+1-0
@@ -118,6 +118,7 @@ __simd_callee__ inline void asc_dump(__ubuf__ T* input, uint32_t desc, uint32_t
118| dump_size | 输入 | 所需要打印的元素数量。 |118| dump_size | 输入 | 所需要打印的元素数量。 |
119 119 
120## 数据类型120## 数据类型
121+ 
121<!-- npu="950" id10 -->122<!-- npu="950" id10 -->
122- 针对Ascend 950PR&950DT系列产品 ：bool、int8_t、uint8_t、hifloat8_t、fp8_e8m0_t、fp8_e5m2_t、fp8_e4m3fn_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float、int64_t、uint64_t。123- 针对Ascend 950PR&950DT系列产品 ：bool、int8_t、uint8_t、hifloat8_t、fp8_e8m0_t、fp8_e5m2_t、fp8_e4m3fn_t、int16_t、uint16_t、half、bfloat16_t、int32_t、uint32_t、float、int64_t、uint64_t。
123<!-- end id10 -->124<!-- end id10 -->
Mdocs/zh/api/appendix/deprecated_interface.md+4-0
@@ -40,6 +40,7 @@
40 此接口后续版本会废弃，请使用[asc_set_atomic_min_int32](../SIMD-API/c_api/atomic/datamove_atomic/asc_set_atomic_min.md)接口。40 此接口后续版本会废弃，请使用[asc_set_atomic_min_int32](../SIMD-API/c_api/atomic/datamove_atomic/asc_set_atomic_min.md)接口。
41 41 
42<!-- npu="950" id2 -->42<!-- npu="950" id2 -->
43+ 
43### 系统变量44### 系统变量
44 45 
45- [asc_get_block_idx](../SIMD-API/c_api/deprecated_interface/asc_get_block_idx_deprecated.md)接口46- [asc_get_block_idx](../SIMD-API/c_api/deprecated_interface/asc_get_block_idx_deprecated.md)接口
@@ -59,6 +60,7 @@
59<!-- end id2 -->60<!-- end id2 -->
60 61 
61<!-- npu="950" id3 -->62<!-- npu="950" id3 -->
63+ 
62### Reg矢量计算64### Reg矢量计算
63 65 
64#### Reg数据搬入66#### Reg数据搬入
@@ -212,6 +214,7 @@
212<!-- end id3 -->214<!-- end id3 -->
213 215 
214<!-- npu="950" id4 -->216<!-- npu="950" id4 -->
217+ 
215#### 逻辑计算218#### 逻辑计算
216 219 
217- [asc_and](../SIMD-API/c_api/deprecated_interface/asc_and_deprecated.md)接口（浮点类型重载）220- [asc_and](../SIMD-API/c_api/deprecated_interface/asc_and_deprecated.md)接口（浮点类型重载）
@@ -508,6 +511,7 @@
508 511 
509 使用`uint8_t unit_flag`入参的`asc_mmad_sparse()`接口已废弃，请使用[asc_mmad_sparse](../SIMD-API/c_api/cube_compute/asc_mmad_sparse.md)接口。`asc_mmad_sparse_sync()`接口已废弃，请使用[asc_mmad_sparse](../SIMD-API/c_api/cube_compute/asc_mmad_sparse.md)接口和[asc_sync](../SIMD-API/c_api/sync/intra_core_sync/asc_sync.md)接口。512 使用`uint8_t unit_flag`入参的`asc_mmad_sparse()`接口已废弃，请使用[asc_mmad_sparse](../SIMD-API/c_api/cube_compute/asc_mmad_sparse.md)接口。`asc_mmad_sparse_sync()`接口已废弃，请使用[asc_mmad_sparse](../SIMD-API/c_api/cube_compute/asc_mmad_sparse.md)接口和[asc_sync](../SIMD-API/c_api/sync/intra_core_sync/asc_sync.md)接口。
510<!-- end id6 -->513<!-- end id6 -->
514+ 
511### 矩阵计算搬运515### 矩阵计算搬运
512 516 
513- [asc_copy_gm2l0a_sync（废弃）](../SIMD-API/c_api/deprecated_interface/asc_copy_gm2l0a_sync_deprecated.md)接口517- [asc_copy_gm2l0a_sync（废弃）](../SIMD-API/c_api/deprecated_interface/asc_copy_gm2l0a_sync_deprecated.md)接口
Mdocs/zh/api/appendix/mmad_mx_interface_boundary_value_summary.md+1-0
@@ -7,6 +7,7 @@
7<!-- end id1 -->7<!-- end id1 -->
8 8 
9## 边界值说明9## 边界值说明
10+ 
10- fp4x2_e1m2_t和fp4x2_e2m1_t数据类型没有+inf/-inf/nan定义，从其他数据类型的+inf/-inf/nan转换为这两种数据类型时会被直接转换为边界值或0。值转换情况如下表所示：11- fp4x2_e1m2_t和fp4x2_e2m1_t数据类型没有+inf/-inf/nan定义，从其他数据类型的+inf/-inf/nan转换为这两种数据类型时会被直接转换为边界值或0。值转换情况如下表所示：
11 12
12 **表1** fp4x2_e1m2_t和fp4x2_e2m1_t数据类型边界值转换结果13 **表1** fp4x2_e1m2_t和fp4x2_e2m1_t数据类型边界值转换结果
Mdocs/zh/guide/getting_started/ascend_c_overview_and_learning_path.md+2-0
@@ -1,6 +1,7 @@
1# Ascend C概述1# Ascend C概述
2 2 
3## CANN概述3## CANN概述
4+ 
4AI（Artificial Intelligence，人工智能）芯片是专为加速人工智能计算任务设计的专用处理器，凭借高度并行化的硬件架构，可高效应对深度学习场景中的大规模张量运算。NPU（Neural-Network Processing Unit，神经网络处理器）作为AI芯片的典型代表，其核心优势在于适配神经网络任务的高效执行；昇腾AI处理器作为专用NPU，内部集成大量高性能计算核心，进一步强化了这一硬件优势，为AI任务的高效运行提供坚实支撑。5AI（Artificial Intelligence，人工智能）芯片是专为加速人工智能计算任务设计的专用处理器，凭借高度并行化的硬件架构，可高效应对深度学习场景中的大规模张量运算。NPU（Neural-Network Processing Unit，神经网络处理器）作为AI芯片的典型代表，其核心优势在于适配神经网络任务的高效执行；昇腾AI处理器作为专用NPU，内部集成大量高性能计算核心，进一步强化了这一硬件优势，为AI任务的高效运行提供坚实支撑。
5 6 
6在神经网络模型中，算子是构成模型的基础计算单元，封装了针对张量数据的特定数学操作逻辑，如卷积、矩阵乘、激活函数等。神经网络模型的推理与训练过程，本质上是一系列算子按照预设逻辑有序执行、协同完成计算任务的过程。为充分释放昇腾NPU的硬件算力，需构建一套能够高效衔接上层AI框架与底层硬件的软件平台，而CANN（Compute Architecture for Neural Networks）正是华为专为昇腾NPU打造的全栈软件解决方案。7在神经网络模型中，算子是构成模型的基础计算单元，封装了针对张量数据的特定数学操作逻辑，如卷积、矩阵乘、激活函数等。神经网络模型的推理与训练过程，本质上是一系列算子按照预设逻辑有序执行、协同完成计算任务的过程。为充分释放昇腾NPU的硬件算力，需构建一套能够高效衔接上层AI框架与底层硬件的软件平台，而CANN（Compute Architecture for Neural Networks）正是华为专为昇腾NPU打造的全栈软件解决方案。
@@ -8,6 +9,7 @@ AI（Artificial Intelligence，人工智能）芯片是专为加速人工智能
8CANN基于分层架构设计，实现了上层应用与底层硬件的无缝衔接。向上，兼容PyTorch等主流AI框架，确保应用层的兼容性与易用性；向下，深度适配昇腾AI处理器，充分发挥硬件的计算潜能，进而构建起CPU与NPU深度融合的异构计算体系。在该体系中，CANN通过Runtime运行时接口驱动CPU与NPU高效协同调度：CPU完成数据准备、任务拆解与调度后，将算子任务下发至NPU加速执行；计算完成后，结果由CPU回收或根据业务需求进行进一步处理。这一协同机制贯穿算子、模型、应用全层级，有效达成计算性能与运行效率的最优平衡，为昇腾AI处理器算力的充分释放提供核心软件支撑。9CANN基于分层架构设计，实现了上层应用与底层硬件的无缝衔接。向上，兼容PyTorch等主流AI框架，确保应用层的兼容性与易用性；向下，深度适配昇腾AI处理器，充分发挥硬件的计算潜能，进而构建起CPU与NPU深度融合的异构计算体系。在该体系中，CANN通过Runtime运行时接口驱动CPU与NPU高效协同调度：CPU完成数据准备、任务拆解与调度后，将算子任务下发至NPU加速执行；计算完成后，结果由CPU回收或根据业务需求进行进一步处理。这一协同机制贯穿算子、模型、应用全层级，有效达成计算性能与运行效率的最优平衡，为昇腾AI处理器算力的充分释放提供核心软件支撑。
9 10 
10## Ascend C简介11## Ascend C简介
12+ 
11大模型时代，算力需求正以前所未有的速度激增。如何开发高性能算子，充分释放硬件每一分潜能，已成为决定AI应用落地的关键。为帮助开发者高效构建可直接运行于NPU的高性能算子，Ascend C应运而生。Ascend C是CANN生态下，面向昇腾AI处理器推出的专用算子开发语言，原生兼容C与C++标准规范。此外，Ascend C秉持**开放芯片完备编程能力，支撑极致性能**的设计理念，构建多层级API体系，满足多场景算子开发诉求，让您能够根据项目需求、团队技能和性能目标，在开发效率与运行性能之间灵活权衡，实现最优平衡。13大模型时代，算力需求正以前所未有的速度激增。如何开发高性能算子，充分释放硬件每一分潜能，已成为决定AI应用落地的关键。为帮助开发者高效构建可直接运行于NPU的高性能算子，Ascend C应运而生。Ascend C是CANN生态下，面向昇腾AI处理器推出的专用算子开发语言，原生兼容C与C++标准规范。此外，Ascend C秉持**开放芯片完备编程能力，支撑极致性能**的设计理念，构建多层级API体系，满足多场景算子开发诉求，让您能够根据项目需求、团队技能和性能目标，在开发效率与运行性能之间灵活权衡，实现最优平衡。
12 14 
13为满足不同场景的诉求，Ascend C构建了从单指令抽象、单核公共算法到多核算子的多层级API体系，其总体逻辑架构如下：15为满足不同场景的诉求，Ascend C构建了从单指令抽象、单核公共算法到多核算子的多层级API体系，其总体逻辑架构如下：
Mdocs/zh/guide/operator_practice/simd_operator_optimization/memory_access/avoid_ub_bank_conflict/avoid_bank_conflict_npu_arch_3510.md+2-0
@@ -42,6 +42,7 @@ UB_ADDR[18:0] = {BANK_DEPTH[8:0], BANK[0], BG[2:0], BANK_WIDTH[4:0]}
42**注：bank冲突的场景与UB的规格密切相关，规格的变化通常会导致bank冲突场景的变化。**42**注：bank冲突的场景与UB的规格密切相关，规格的变化通常会导致bank冲突场景的变化。**
43 43 
44### 基于连续访存的bank冲突分析44### 基于连续访存的bank冲突分析
45+ 
45以下对连续访存的冲突类型进行解读：46以下对连续访存的冲突类型进行解读：
46 47 
47- 读写冲突48- 读写冲突
@@ -239,6 +240,7 @@ UB_ADDR[18:0] = {BANK_DEPTH[8:0], BANK[0], BG[2:0], BANK_WIDTH[4:0]}
239 </table>240 </table>
240 241 
241### 基于离散访存的bank冲突分析242### 基于离散访存的bank冲突分析
243+ 
242以下对离散访存下的冲突进行解读：244以下对离散访存下的冲突进行解读：
243 245 
244- 当gather/scatter访存数据分布在不同的subbank中时，不发生冲突。246- 当gather/scatter访存数据分布在不同的subbank中时，不发生冲突。
Mdocs/zh/guide/operator_practice/simd_operator_optimization/vector_compute/vf_optimization/dual_issue_optimization.md+2-0
@@ -91,6 +91,7 @@ for (uint16_t i = 0; i < 32; ++i) { // VLoop-2
91> 循环展开通过消除指令依赖，提升指令双发能力来优化VF性能，需根据实际计算过程使用，不保证产生收益。展开时，若寄存器超出数量上限，可能会造成性能劣化或其他非预期行为。91> 循环展开通过消除指令依赖，提升指令双发能力来优化VF性能，需根据实际计算过程使用，不保证产生收益。展开时，若寄存器超出数量上限，可能会造成性能劣化或其他非预期行为。
92 92 
93### 手动展开VF循环93### 手动展开VF循环
94+ 
94```cpp95```cpp
95for (uint16_t i = 0; i < 32; ++i) { 96for (uint16_t i = 0; i < 32; ++i) {
96// for循环间没有依赖关系：i=0,i=1可以并行执行，但是由于循环内数据依赖指令过多导致i=1的指令无法加载进执行队列中97// for循环间没有依赖关系：i=0,i=1可以并行执行，但是由于循环内数据依赖指令过多导致i=1的指令无法加载进执行队列中
@@ -127,6 +128,7 @@ for (uint16_t i = 0; i < 8; ++i) { // 32做4展开
127```128```
128 129 
129### 编译器自动展开130### 编译器自动展开
131+ 
130```cpp132```cpp
131#pragma unroll 4 133#pragma unroll 4
132for (uint16_t i = 0; i < 32; ++i) { 134for (uint16_t i = 0; i < 32; ++i) {
Mdocs/zh/guide/operator_practice/simt_operator_optimization/execution_config/thread_block_config.md+1-0
@@ -19,6 +19,7 @@
19本优化方法对应的样例请参见[核函数（Kernel）的启动配置性能调优样例](../../../../../../examples/03_simt_api/03_best_practices/01_execution_conf_optimizations/grid_dim_config/README.md)。19本优化方法对应的样例请参见[核函数（Kernel）的启动配置性能调优样例](../../../../../../examples/03_simt_api/03_best_practices/01_execution_conf_optimizations/grid_dim_config/README.md)。
20 20 
21## 大数据量场景21## 大数据量场景
22+ 
22下面以Gather算子为例，对比大数据量场景（`1024*2048`，即2097152个元素）下不同线程块数量配置的性能差异。23下面以Gather算子为例，对比大数据量场景（`1024*2048`，即2097152个元素）下不同线程块数量配置的性能差异。
23 24 
24【反例】一个线程处理一个元素，随数据量增加线程块个数，导致需要启动和调度的线程块大幅增加。25【反例】一个线程处理一个元素，随数据量增加线程块个数，导致需要启动和调度的线程块大幅增加。
Mdocs/zh/guide/operator_practice/typical_operator_cases/simd_operator_impl/add_vector_operator/getting_started.md+1-0
@@ -125,6 +125,7 @@ AscendC::DumpTensor(zLocal, 1, 32);
125## 性能调试125## 性能调试
126 126 
127### msOpProf工具介绍127### msOpProf工具介绍
128+ 
128msOpProf工具是单算子性能分析工具。包含msopprof和msopprof simulator两种使用方式。该工具协助用户定位算子内存、算子代码以及算子指令的异常，实现全方位的算子调优。当前支持基于不同运行模式（上板或仿真）和不同文件形式（可执行文件或算子二进制.o文件）进行性能数据的采集和自动解析。129msOpProf工具是单算子性能分析工具。包含msopprof和msopprof simulator两种使用方式。该工具协助用户定位算子内存、算子代码以及算子指令的异常，实现全方位的算子调优。当前支持基于不同运行模式（上板或仿真）和不同文件形式（可执行文件或算子二进制.o文件）进行性能数据的采集和自动解析。
129 130 
130- 上板性能采集131- 上板性能采集
Mdocs/zh/guide/programming_guide/advanced_programming/aclnn_operator_development/appendix/cmake_function_reference.md+2-0
@@ -215,6 +215,7 @@ npu_op_kernel_options(<target_name> <op_type> [COMPUTE_UNIT <soc_version>... | S
215npu_op_kernel_options(ascendc_kernels ALL OPTIONS --save-temp-files -g) #为算子添加编译选项215npu_op_kernel_options(ascendc_kernels ALL OPTIONS --save-temp-files -g) #为算子添加编译选项
216```216```
217npu_op_kernel_options(ascendc_kernels AddCustom SOC_SERIES Ascendxxx Ascendyyy OPTIONS -DASCENDC_DEBUG)217npu_op_kernel_options(ascendc_kernels AddCustom SOC_SERIES Ascendxxx Ascendyyy OPTIONS -DASCENDC_DEBUG)
218+ 
218### npu_op_kernel_sources219### npu_op_kernel_sources
219 220 
220描述核函数（Kernel）目标的源码信息，包括设置算子的核函数（Kernel）实现文件和源码路径等。221描述核函数（Kernel）目标的源码信息，包括设置算子的核函数（Kernel）实现文件和源码路径等。
@@ -251,6 +252,7 @@ npu_op_kernel_sources(ascendc_kernels
251 KERNEL_FILE add_custom.cpp252 KERNEL_FILE add_custom.cpp
252)253)
253```254```
255+ 
254### npu_op_device_tiling_library256### npu_op_device_tiling_library
255 257 
256创建Device侧Tiling库。使用该选项时，package的类型仅支持配置为RUN（run包模式）。258创建Device侧Tiling库。使用该选项时，package的类型仅支持配置为RUN（run包模式）。
Mdocs/zh/guide/programming_guide/advanced_programming/aclnn_operator_development/compilation_and_deployment/basic_process.md+5-0
@@ -148,6 +148,7 @@ export ASCENDC_BUILD_LOG_DIR=/home/build_log/
148```148```
149 149 
150<a id="算子包部署"></a>150<a id="算子包部署"></a>
151+ 
151## 算子包部署<a id="operator-package-deployment"></a>152## 算子包部署<a id="operator-package-deployment"></a>
152 153 
153编译生成的算子安装包需要部署到运行环境。154编译生成的算子安装包需要部署到运行环境。
@@ -184,6 +185,7 @@ source /usr/local/Ascend/cann/set_env.sh
184> - 算子包要求架构与系统架构一致，且满足glibc版本依赖，否则可使用`--force`强制安装。185> - 算子包要求架构与系统架构一致，且满足glibc版本依赖，否则可使用`--force`强制安装。
185 186 
186<a id="部署后的目录结构"></a>187<a id="部署后的目录结构"></a>
188+ 
187### 部署后的目录结构189### 部署后的目录结构
188 190 
189以默认安装场景为例：191以默认安装场景为例：
@@ -223,6 +225,7 @@ source /usr/local/Ascend/cann/set_env.sh
223```225```
224 226 
225<a id="配置自定义算子优先级"></a>227<a id="配置自定义算子优先级"></a>
228+ 
226### 配置自定义算子优先级229### 配置自定义算子优先级
227 230 
228多个算子包共存时，如果包含相同OpType的算子，优先级高的算子包中的版本生效。231多个算子包共存时，如果包含相同OpType的算子，优先级高的算子包中的版本生效。
@@ -313,6 +316,7 @@ source /usr/local/Ascend/cann/set_env.sh
313---316---
314 317 
315<a id="编译组织"></a>318<a id="编译组织"></a>
319+ 
316## 编译组织（进阶）<a id="advanced-build-organization"></a>320## 编译组织（进阶）<a id="advanced-build-organization"></a>
317 321 
318算子工程通过CMake组织编译。如需自行组织工程或用CANN提供的CMake modules定制编译流程，可参考以下内容。322算子工程通过CMake组织编译。如需自行组织工程或用CANN提供的CMake modules定制编译流程，可参考以下内容。
@@ -613,6 +617,7 @@ npu_op_package_add(${package_name}
613完成编译组织之后，可按照[基本流程](#基本流程)的三个步骤完成算子编译与部署。617完成编译组织之后，可按照[基本流程](#基本流程)的三个步骤完成算子编译与部署。
614 618 
615<a id="自定义编译选项"></a>619<a id="自定义编译选项"></a>
620+ 
616## 自定义编译选项（进阶）<a id="custom-compile-options"></a>621## 自定义编译选项（进阶）<a id="custom-compile-options"></a>
617 622 
618在CMakeLists.txt中，用`npu_op_kernel_options`为核函数（Kernel）侧代码添加编译选项：623在CMakeLists.txt中，用`npu_op_kernel_options`为核函数（Kernel）侧代码添加编译选项：
Mdocs/zh/guide/programming_guide/advanced_programming/aclnn_operator_development/compilation_and_deployment/compilation_acceleration.md+2-0
@@ -147,6 +147,7 @@ ccache --show-stats -v
147 147 
148 148 
149<a id="选择性编译"></a>149<a id="选择性编译"></a>
150+ 
150## 选择性编译151## 选择性编译
151 152 
152选择性编译的核心思路是：开发调试阶段只编译当前真正需要验证的部分，减少无关核函数（Kernel）变体、无关芯片和无关产物带来的编译开销。常用策略如下：153选择性编译的核心思路是：开发调试阶段只编译当前真正需要验证的部分，减少无关核函数（Kernel）变体、无关芯片和无关产物带来的编译开销。常用策略如下：
@@ -245,6 +246,7 @@ npu_op_kernel_options(ascendc_kernels AddCustomTemplate COMPUTE_UNIT Ascendxxyy
245cmake -S . -B build_out --preset=default -DASCEND_SOC_SERIES=ascendxxx246cmake -S . -B build_out --preset=default -DASCEND_SOC_SERIES=ascendxxx
246```247```
247也可以使用参数`-DASCEND_COMPUTE_UNIT=ascendxxxyy`；两个参数不能同时传入。248也可以使用参数`-DASCEND_COMPUTE_UNIT=ascendxxxyy`；两个参数不能同时传入。
249+ 
248### 跳过核函数（Kernel）二进制编译250### 跳过核函数（Kernel）二进制编译
249 251 
250如果当前只需要生成源码包，或暂时不需要核函数（Kernel） `.o`二进制，可关闭二进制包生成。顶层`CMakeLists.txt`中应使用变量形式传递`ENABLE_BINARY_PACKAGE`：252如果当前只需要生成源码包，或暂时不需要核函数（Kernel） `.o`二进制，可关闭二进制包生成。顶层`CMakeLists.txt`中应使用变量形式传递`ENABLE_BINARY_PACKAGE`：
Mdocs/zh/guide/programming_guide/advanced_programming/aclnn_operator_development/invocation/single_operator_api_call.md+1-0
@@ -313,6 +313,7 @@ aclnnStatus aclnnXxxTensorGetWorkspaceSize(const aclTensor *src, ..., aclTensor
313aclnnStatus aclnnXxxGetWorkspaceSize(const aclIntArray *x0, const aclBoolArray *x1, const aclFloatArray *x2, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor);313aclnnStatus aclnnXxxGetWorkspaceSize(const aclIntArray *x0, const aclBoolArray *x1, const aclFloatArray *x2, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor);
314aclnnStatus aclnnXxxTensorGetWorkspaceSize(const aclTensor *x0, const aclTensor *x1, const aclTensor *x2, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor);314aclnnStatus aclnnXxxTensorGetWorkspaceSize(const aclTensor *x0, const aclTensor *x1, const aclTensor *x2, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor);
315```315```
316+ 
316## 相关文档317## 相关文档
317 318 
318- [快速入门](../aclnn_quick_start.md) — 创建并跑通AddCustom算子工程。319- [快速入门](../aclnn_quick_start.md) — 创建并跑通AddCustom算子工程。
Mdocs/zh/guide/programming_guide/advanced_programming/memory_model/cache_coherence.md+4-0
@@ -220,6 +220,7 @@ T6 Scalar.LOAD A -> 1
220还有一类问题来自Cache Line粒度。普通Scalar写后续以64B Cache Line为单位写回；多个核即使写的是同一条Cache Line中的不同字段，也可能在各自写回整条Cache Line时覆盖对方结果。此类场景可以通过让不同核操作的地址至少相隔Cache Line大小，或使用ReadGmBypassDCache和WriteGmBypassDCache降低风险。220还有一类问题来自Cache Line粒度。普通Scalar写后续以64B Cache Line为单位写回；多个核即使写的是同一条Cache Line中的不同字段，也可能在各自写回整条Cache Line时覆盖对方结果。此类场景可以通过让不同核操作的地址至少相隔Cache Line大小，或使用ReadGmBypassDCache和WriteGmBypassDCache降低风险。
221 221 
222<!-- npu="950" id3 -->222<!-- npu="950" id3 -->
223+ 
223#### Atomic访问导致的DCache缓存一致性224#### Atomic访问导致的DCache缓存一致性
224 225 
225[NPU架构版本3510](../../language_extension/simd_builtin_keywords.md)支持Scalar原子操作，Scalar原子操作可以看成对某个GM地址执行的读-改-写事件，但是Scalar原子操作会绕过DCache，因此会导致缓存一致性。例如C API [asc_atomic_add](../../../../api/SIMD-API/c_api/atomic/scalar_atomic/asc_atomic_add.md)，该接口在指定GM地址上执行原子加操作。原子操作保证的是这一次读-改-写不会被其他原子更新打断，但它不会自动清理同一地址在DCache中的旧副本或Dirty副本。226[NPU架构版本3510](../../language_extension/simd_builtin_keywords.md)支持Scalar原子操作，Scalar原子操作可以看成对某个GM地址执行的读-改-写事件，但是Scalar原子操作会绕过DCache，因此会导致缓存一致性。例如C API [asc_atomic_add](../../../../api/SIMD-API/c_api/atomic/scalar_atomic/asc_atomic_add.md)，该接口在指定GM地址上执行原子加操作。原子操作保证的是这一次读-改-写不会被其他原子更新打断，但它不会自动清理同一地址在DCache中的旧副本或Dirty副本。
@@ -448,6 +449,7 @@ ICache用于缓存Scalar单元最近使用或频繁使用的指令。Scalar读
448从缓存一致性角度看，ICache缓存的是指令而不是普通数据，并且是只读Cache，因此一般不涉及普通数据路径上的Cache副本不一致问题。449从缓存一致性角度看，ICache缓存的是指令而不是普通数据，并且是只读Cache，因此一般不涉及普通数据路径上的Cache副本不一致问题。
449 450 
450<!-- npu="950" id4 -->451<!-- npu="950" id4 -->
452+ 
451### SIMT DCache与GM的缓存一致性453### SIMT DCache与GM的缓存一致性
452 454 
453#### SIMT DCache访问路径455#### SIMT DCache访问路径
@@ -608,6 +610,7 @@ C API中提供了[`asc_ub_dcci_single`](../../../../api/SIMD-API/c_api/cache_ctr
608下面按atomic事件发起路径分类说明：<!-- npu="950" id6 -->Scalar原子、<!-- end id6 -->MTE3随路原子、FixPipe随路原子<!-- npu="950" id7 -->和SIMT原子<!-- end id7 -->。本节只说明这些atomic事件放入一致性分析时需要关注的Cache行为，接口细节可参考对应API文档。610下面按atomic事件发起路径分类说明：<!-- npu="950" id6 -->Scalar原子、<!-- end id6 -->MTE3随路原子、FixPipe随路原子<!-- npu="950" id7 -->和SIMT原子<!-- end id7 -->。本节只说明这些atomic事件放入一致性分析时需要关注的Cache行为，接口细节可参考对应API文档。
609 611 
610<!-- npu="950" id8 -->612<!-- npu="950" id8 -->
613+ 
611### Scalar原子操作与DCache一致性614### Scalar原子操作与DCache一致性
612 615 
613[NPU架构版本3510](../../language_extension/simd_builtin_keywords.md)支持Scalar原子操作，Scalar原子操作和普通Scalar读写GM混用时，需要按DCache与GM的一致性来分析，详细处理方式参考[Atomic访问导致的DCache缓存一致性](#atomic访问导致的dcache缓存一致性)。616[NPU架构版本3510](../../language_extension/simd_builtin_keywords.md)支持Scalar原子操作，Scalar原子操作和普通Scalar读写GM混用时，需要按DCache与GM的一致性来分析，详细处理方式参考[Atomic访问导致的DCache缓存一致性](#atomic访问导致的dcache缓存一致性)。
@@ -635,6 +638,7 @@ FixPipe常用于`L0C Buffer -> GM`或相关矩阵搬出路径。对一致性分
635- 如果多个流水或多个核都对同一GM地址做随路原子，原子操作只保证每次GM读-改-写不被打断；不同流水、不同核之间的启动顺序仍应由同步或算法协议保证。638- 如果多个流水或多个核都对同一GM地址做随路原子，原子操作只保证每次GM读-改-写不被打断；不同流水、不同核之间的启动顺序仍应由同步或算法协议保证。
636 639 
637<!-- npu="950" id9 -->640<!-- npu="950" id9 -->
641+ 
638### SIMT原子操作与缓存一致性642### SIMT原子操作与缓存一致性
639 643 
640[NPU架构版本3510](../../language_extension/simd_builtin_keywords.md)支持SIMT原子操作，SIMT原子操作可能作用于UB，也可能作用于GM；接口例子可参考[asc_atomic_add](../../../../api/SIMT-API/atomic_operations/asc_atomic_add.md)。它保证的是同一地址上这一次读-改-写具有原子性，不保证多个线程之间的执行顺序。644[NPU架构版本3510](../../language_extension/simd_builtin_keywords.md)支持SIMT原子操作，SIMT原子操作可能作用于UB，也可能作用于GM；接口例子可参考[asc_atomic_add](../../../../api/SIMT-API/atomic_operations/asc_atomic_add.md)。它保证的是同一地址上这一次读-改-写具有原子性，不保证多个线程之间的执行顺序。
Mdocs/zh/guide/programming_guide/advanced_programming/memory_model/memory_consistency.md+2-0
@@ -115,6 +115,7 @@ Scalar负责标量控制、地址计算、标量计算、Scalar访存以及发
115MTE/Fixpipe指令无论是否访问同地址，指令间内存访问的实际完成顺序或对外可见顺序可能与代码顺序不一致，硬件不保证内存一致性。如果前后指令存在重叠地址或数据依赖，需要显式等待前序指令的数据读写完成。115MTE/Fixpipe指令无论是否访问同地址，指令间内存访问的实际完成顺序或对外可见顺序可能与代码顺序不一致，硬件不保证内存一致性。如果前后指令存在重叠地址或数据依赖，需要显式等待前序指令的数据读写完成。
116 116 
117<!-- npu="950" id4 -->117<!-- npu="950" id4 -->
118+ 
118##### SIMT流水119##### SIMT流水
119 120 
120SIMT流水用于SIMT编程模型中的线程级执行，适合复杂控制流和不规则访存场景。SIMT线程可以访问GM，也支持线程同步、MEMBAR和原子操作。121SIMT流水用于SIMT编程模型中的线程级执行，适合复杂控制流和不规则访存场景。SIMT线程可以访问GM，也支持线程同步、MEMBAR和原子操作。
@@ -503,6 +504,7 @@ T3 sync S->MTE2
503`DSB`和流水同步只处理访问完成顺序和执行先后关系，保证内存一致性，不负责写回或失效DCache副本。如果Scalar写GM经过DCache，或者后续读取路径可能命中旧Cache副本，还需要参考[缓存一致性](cache_coherence.md)处理DCache相关问题。504`DSB`和流水同步只处理访问完成顺序和执行先后关系，保证内存一致性，不负责写回或失效DCache副本。如果Scalar写GM经过DCache，或者后续读取路径可能命中旧Cache副本，还需要参考[缓存一致性](cache_coherence.md)处理DCache相关问题。
504 505 
505<!-- npu="950" id11 -->506<!-- npu="950" id11 -->
507+ 
506### SIMT内存一致性508### SIMT内存一致性
507 509 
508SIMT内存一致性需要区分线程内和线程间：510SIMT内存一致性需要区分线程内和线程间：
Mdocs/zh/guide/programming_guide/advanced_programming/super_kernel/operator_adaptation.md+1-0
@@ -52,6 +52,7 @@ Cache刷新机制示意图如下图所示：
52| [AscendC::GetBlockNum()](../../../../api/SIMD-API/basic_api/tool_interface/system_resources_and_variables/GetBlockNum.md) | 获取当前任务配置的核数 |52| [AscendC::GetBlockNum()](../../../../api/SIMD-API/basic_api/tool_interface/system_resources_and_variables/GetBlockNum.md) | 获取当前任务配置的核数 |
53 53 
54<!-- npu="A3,910b" id1 -->54<!-- npu="A3,910b" id1 -->
55+ 
55## TPipe析构约束56## TPipe析构约束
56 57 
57针对Atlas A2系列产品、Atlas A3系列产品，在GE入图及npugraph_ex后端aclnn调用场景，子算子TPipe对象析构（[TPipe::Destroy](../../../../api/SIMD-API/basic_api/resource_management/TPipe/Destroy.md)）接口内部的`AscendC::PipeBarrier<PIPE_ALL>()`指令会被去除，如果算子内部使用了多个TPipe对象或手动调用Destroy函数时开发者需自行保障TPipe对象间流水的同步。58针对Atlas A2系列产品、Atlas A3系列产品，在GE入图及npugraph_ex后端aclnn调用场景，子算子TPipe对象析构（[TPipe::Destroy](../../../../api/SIMD-API/basic_api/resource_management/TPipe/Destroy.md)）接口内部的`AscendC::PipeBarrier<PIPE_ALL>()`指令会被去除，如果算子内部使用了多个TPipe对象或手动调用Destroy函数时开发者需自行保障TPipe对象间流水的同步。
Mdocs/zh/guide/programming_guide/compilation_and_execution/async_execution.md+6-0
@@ -1,4 +1,5 @@
1# 异步执行1# 异步执行
2+ 
2当开发者完成一个核函数（Kernel）的编写后，需要通过特定的下发/启动语法将其提交到指定的Stream（计算流）上执行。例如：3当开发者完成一个核函数（Kernel）的编写后，需要通过特定的下发/启动语法将其提交到指定的Stream（计算流）上执行。例如：
3```c++4```c++
4 kernel<<<numBlocks, ubufDynSize, stream>>>(参数列表);5 kernel<<<numBlocks, ubufDynSize, stream>>>(参数列表);
@@ -22,6 +23,7 @@
22在这个过程中我们提到了Stream创建与销毁，内存申请与拷贝，异步调用与同步等待。在本节的其余部分，我们将解释这些CANN Runtime的元素和对应的API。23在这个过程中我们提到了Stream创建与销毁，内存申请与拷贝，异步调用与同步等待。在本节的其余部分，我们将解释这些CANN Runtime的元素和对应的API。
23 24 
24## Stream管理25## Stream管理
26+ 
25Stream（计算流）本质上是一个有序的任务队列，为开发者提供了一种表示“先序执行”的抽象机制，程序可向队列中提交各类操作，例如内存拷贝、核函数（Kernel）下发等，所有任务将按入队顺序依次执行。单个计算流内，队列头部任务优先执行，执行完成后自动出队；后续队列任务依次前移，等待调度执行。同一个计算流中的任务具备串行执行特性，严格遵循入队先后顺序运行。27Stream（计算流）本质上是一个有序的任务队列，为开发者提供了一种表示“先序执行”的抽象机制，程序可向队列中提交各类操作，例如内存拷贝、核函数（Kernel）下发等，所有任务将按入队顺序依次执行。单个计算流内，队列头部任务优先执行，执行完成后自动出队；后续队列任务依次前移，等待调度执行。同一个计算流中的任务具备串行执行特性，严格遵循入队先后顺序运行。
26 28 
27应用程序可同时创建并使用多个计算流。多流场景下，运行时会根据AI处理器的硬件资源占用状态，从存有待处理任务的计算流中择优调度任务执行。开发者可为计算流配置优先级，该优先级将作为运行时调度的参考依据，用以调整任务调度策略，但无法强制保证固定的执行顺序。29应用程序可同时创建并使用多个计算流。多流场景下，运行时会根据AI处理器的硬件资源占用状态，从存有待处理任务的计算流中择优调度任务执行。开发者可为计算流配置优先级，该优先级将作为运行时调度的参考依据，用以调整任务调度策略，但无法强制保证固定的执行顺序。
@@ -37,6 +39,7 @@ CANN Runtime内置默认计算流，所有未显式指定计算流的运算操
37- **指定Stream同步**：`aclrtSynchronizeStream(stream)`。39- **指定Stream同步**：`aclrtSynchronizeStream(stream)`。
38 40 
39## Memory（内存）管理41## Memory（内存）管理
42+ 
40在异构计算架构中，系统由主机（Host）和设备（Device）组成。Host和Device各自拥有独立的内存，Host内存是指AI处理器所在服务器的主机内存（即CPU内存），而Device内存则是指AI处理器自带的设备内存。43在异构计算架构中，系统由主机（Host）和设备（Device）组成。Host和Device各自拥有独立的内存，Host内存是指AI处理器所在服务器的主机内存（即CPU内存），而Device内存则是指AI处理器自带的设备内存。
41 44 
42内存管理中要做好的两件事是：45内存管理中要做好的两件事是：
@@ -44,6 +47,7 @@ CANN Runtime内置默认计算流，所有未显式指定计算流的运算操
442. **高效访问内存**：算子在Device上执行过程中，访问Device上的Device内存数据性能更高。为此，Runtime提供了Host与Device之间互相拷贝内存的接口，支持同步和异步方式，例如aclrtMemcpy和aclrtMemcpyAsync等，以便开发者更好地规划数据的存储与访问。472. **高效访问内存**：算子在Device上执行过程中，访问Device上的Device内存数据性能更高。为此，Runtime提供了Host与Device之间互相拷贝内存的接口，支持同步和异步方式，例如aclrtMemcpy和aclrtMemcpyAsync等，以便开发者更好地规划数据的存储与访问。
45 48 
46## 异步调用49## 异步调用
50+ 
47为了能够更好的发挥CPU和AI处理器的硬件利用率，CANN Runtime遵循“下发即返回”的异步执行策略，使Host在提交H2D/D2H拷贝或核函数（Kernel）计算任务后立即继续执行后续逻辑，从而让主机侧调度、设备侧计算、主机与设备之间的数据传输尽可能重叠执行，隐藏传输和调度开销，提高端到端吞吐率。51为了能够更好的发挥CPU和AI处理器的硬件利用率，CANN Runtime遵循“下发即返回”的异步执行策略，使Host在提交H2D/D2H拷贝或核函数（Kernel）计算任务后立即继续执行后续逻辑，从而让主机侧调度、设备侧计算、主机与设备之间的数据传输尽可能重叠执行，隐藏传输和调度开销，提高端到端吞吐率。
48 52 
49典型<<<\>\>\>异步调用的场景是：将主机与设备计算操作重叠执行，从而提升主机侧与设备侧计算资源利用率。53典型<<<\>\>\>异步调用的场景是：将主机与设备计算操作重叠执行，从而提升主机侧与设备侧计算资源利用率。
@@ -56,6 +60,7 @@ CANN Runtime内置默认计算流，所有未显式指定计算流的运算操
56 60 
57 61 
58## 同步等待62## 同步等待
63+ 
59核函数（Kernel）的调用是异步的，在核函数（Kernel）调用结束之后，控制权立刻返回给主机端，需要在host等待核函数（Kernel）的执行完成，可以通过以下接口强制Host程序等待核函数（Kernel）执行完毕：64核函数（Kernel）的调用是异步的，在核函数（Kernel）调用结束之后，控制权立刻返回给主机端，需要在host等待核函数（Kernel）的执行完成，可以通过以下接口强制Host程序等待核函数（Kernel）执行完毕：
60|同步接口|接口描述|65|同步接口|接口描述|
61|-|-|66|-|-|
@@ -63,4 +68,5 @@ CANN Runtime内置默认计算流，所有未显式指定计算流的运算操
63|`aclrtSynchronizeStream()` |阻塞Host侧当前线程直到指定Stream中的所有操作都已完成。 |68|`aclrtSynchronizeStream()` |阻塞Host侧当前线程直到指定Stream中的所有操作都已完成。 |
64 69 
65## 更多的Runtime高级特性70## 更多的Runtime高级特性
71+ 
66更多的Runtime编程方式及接口，可以参考[《应用开发（C&C++）》](https://hiascend.com/document/redirect/CannCommunityadev)。72更多的Runtime编程方式及接口，可以参考[《应用开发（C&C++）》](https://hiascend.com/document/redirect/CannCommunityadev)。
Mdocs/zh/guide/programming_guide/compilation_and_execution/operator_compilation/ai_core_operator_compilation.md+3-0
@@ -1,7 +1,9 @@
1# AI Core编译基本用法<a name="section229217121411"></a>1# AI Core编译基本用法<a name="section229217121411"></a>
2+ 
2当开发者完成一个核函数（Kernel）的编写，并在Host侧通过<<<\>\>\>方式完成核函数（Kernel）的启动调用后，就需要进行算子源码编译与算子运行，接下来我们介绍通过bisheng编译器完成算子编译的基本用法。3当开发者完成一个核函数（Kernel）的编写，并在Host侧通过<<<\>\>\>方式完成核函数（Kernel）的启动调用后，就需要进行算子源码编译与算子运行，接下来我们介绍通过bisheng编译器完成算子编译的基本用法。
3 4 
4## AI Core SIMD编译5## AI Core SIMD编译
6+ 
5使用bisheng编译Ascend C源文件的基本命令如下，其中--npu-arch=dav-\<npu architecture\>用于指定AI处理器架构版本：7使用bisheng编译Ascend C源文件的基本命令如下，其中--npu-arch=dav-\<npu architecture\>用于指定AI处理器架构版本：
6```shell8```shell
7 bisheng <source_file>.asc -o <output_file> --npu-arch=dav-<npu architecture>9 bisheng <source_file>.asc -o <output_file> --npu-arch=dav-<npu architecture>
@@ -336,6 +338,7 @@ add_executable(demo
336| libascendcl.so | acl相关接口库。 |338| libascendcl.so | acl相关接口库。 |
337 339 
338### 高阶API常用链接库340### 高阶API常用链接库
341+ 
339在使用高阶API时，必须链接以下库，因为这些库是高阶API功能所依赖的。在其他场景下，可以根据具体需求选择是否链接这些库。342在使用高阶API时，必须链接以下库，因为这些库是高阶API功能所依赖的。在其他场景下，可以根据具体需求选择是否链接这些库。
340 343 
341|链接库名称|作用描述|使用场景|动态库路径|344|链接库名称|作用描述|使用场景|动态库路径|
Mdocs/zh/guide/programming_guide/debug_and_tuning/functional_debug/npu_board_debug.md+8-0
@@ -90,11 +90,14 @@ msDebug是一款面向昇腾设备的算子调试工具，用于调试NPU侧运
90### msdebug调试进程90### msdebug调试进程
91 91 
92#### 调试程序编译92#### 调试程序编译
93+ 
93```shell94```shell
94// bisheng编译时增加debug命令行95// bisheng编译时增加debug命令行
95bisheng add_custom.asc -o add_custom -O0 -g96bisheng add_custom.asc -o add_custom -O0 -g
96```97```
98+ 
97#### 调试程序运行99#### 调试程序运行
100+ 
98```shell101```shell
99// 使用msdebug进行调试102// 使用msdebug进行调试
100$ msdebug ./add_custom103$ msdebug ./add_custom
@@ -110,12 +113,14 @@ Breakpoint 1: where = device_debugdata`_ZN17MatmulLeakyKernelIDhDhffE7CopyOutEj_
110```113```
111 114 
112#### 断点设置115#### 断点设置
116+ 
113```shell117```shell
114(msdebug) b matmul_leakyrelu_kernel.cpp:114118(msdebug) b matmul_leakyrelu_kernel.cpp:114
115Breakpoint 1: where = device_debugdata`_ZN17MatmulLeakyKernelIDhDhffE7CopyOutEj_mix_aiv + 240 at matmul_leakyrelu_kernel.cpp:114:14, address = 0x000000000000ff88119Breakpoint 1: where = device_debugdata`_ZN17MatmulLeakyKernelIDhDhffE7CopyOutEj_mix_aiv + 240 at matmul_leakyrelu_kernel.cpp:114:14, address = 0x000000000000ff88
116```120```
117 121 
118#### 内存与变量打印122#### 内存与变量打印
123+ 
119```shell124```shell
120# 打印LocalTensor125# 打印LocalTensor
121(msdebug) p reluOutLocal126(msdebug) p reluOutLocal
@@ -134,6 +139,7 @@ Breakpoint 1: where = device_debugdata`_ZN17MatmulLeakyKernelIDhDhffE7CopyOutEj_
134```139```
135 140 
136#### 单步调试141#### 单步调试
142+ 
137```shell143```shell
138(msdebug) s144(msdebug) s
139Process 177943 stopped145Process 177943 stopped
@@ -148,6 +154,7 @@ Process 177943 stopped
148 ```154 ```
149 155 
150#### 调试信息展示156#### 调试信息展示
157+ 
151 ```shell158 ```shell
152 (msdebug) ascend info cores159 (msdebug) ascend info cores
153 CoreId Type Device Stream Task Block PC stop reason160 CoreId Type Device Stream Task Block PC stop reason
@@ -157,6 +164,7 @@ Process 177943 stopped
157 ```164 ```
158 165 
159### 解析异常算子dump文件166### 解析异常算子dump文件
167+ 
160通过环境变量开启算子异常时Dump出Core文件。168通过环境变量开启算子异常时Dump出Core文件。
161```shell169```shell
162# aic_err_detail_dump: 导出AI Core的内存存储、寄存器和调用栈信息。170# aic_err_detail_dump: 导出AI Core的内存存储、寄存器和调用栈信息。
Mdocs/zh/guide/programming_guide/language_extension/simd_simt_hybrid_builtin_keywords.md+1-0
@@ -123,6 +123,7 @@ kernel_name<<<block_num, dyn_ub_size, stream>>>(args...);
123## VF函数的asc\_vf\_call调用123## VF函数的asc\_vf\_call调用
124 124 
125### SIMT VF的asc\_vf\_call调用125### SIMT VF的asc\_vf\_call调用
126+ 
126SIMD与SIMT混合编程以SIMD核函数（Kernel）作为Device侧入口，在核函数（Kernel）或`__aicore__`函数中通过`asc_vf_call`启动SIMT VF子任务，通过参数配置，启动指定数目的线程，执行指定的SIMT VF函数。其函数原型如下：127SIMD与SIMT混合编程以SIMD核函数（Kernel）作为Device侧入口，在核函数（Kernel）或`__aicore__`函数中通过`asc_vf_call`启动SIMT VF子任务，通过参数配置，启动指定数目的线程，执行指定的SIMT VF函数。其函数原型如下：
127 128 
128```cpp129```cpp
Mdocs/zh/guide/programming_guide/programming_model/ai_core_simd_programming/c_pointer_programming/cube_matrix_computation.md+1-0
@@ -63,6 +63,7 @@ Cube单元采用分块计算逻辑，硬件最小计算粒度为分形块，可
63ND格式是通用的N维张量格式，数据在内存中连续线性存放，通常存在于`Global Memory (GM)`中。在进行矩阵运算前，通常需要将其转换为Nz格式以适配Cube计算单元。63ND格式是通用的N维张量格式，数据在内存中连续线性存放，通常存在于`Global Memory (GM)`中。在进行矩阵运算前，通常需要将其转换为Nz格式以适配Cube计算单元。
64 64 
65#### Nz 格式65#### Nz 格式
66+ 
66Nz格式主要用于L1 Buffer和L0C Buffer中存放数据。针对[NPU架构版本3510](../../../language_extension/simd_builtin_keywords.md)，L0A Buffer也采用Nz分形格式。该格式采用大N（外部列主序）+ 小z（内部行主序）的方式排布。67Nz格式主要用于L1 Buffer和L0C Buffer中存放数据。针对[NPU架构版本3510](../../../language_extension/simd_builtin_keywords.md)，L0A Buffer也采用Nz分形格式。该格式采用大N（外部列主序）+ 小z（内部行主序）的方式排布。
67 68 
68<img src="../../../../figures/nz_format.png" alt="Nz数据排布格式" width="400"/>69<img src="../../../../figures/nz_format.png" alt="Nz数据排布格式" width="400"/>