Star1189
1200
Code
Code
Issues140
Pull Requests415
Actions
Discussion
Wiki
Analysis
Star1189
1200
Merged
docs: add blank lines around code fences in docs/zh to meet format rules #6518
ai_xinAdd on Sep 30
docs: add blank lines around code fences in docs/zh to meet format rules #6518
Merged
ai_xinAdd on Sep 30
Merged fromai_xin/asc-devkitmasterintocann/asc-devkitmaster
Total: 132 files changed+224-26
Mdocs/zh/api/SIMD-API/adv_api/cube_compute/Matmul_Kernel/Matmul_template_params.md+1-0
@@ -81,6 +81,7 @@
81## 函数原型81## 函数原型
82 82 
83Matmul模板参数如下：83Matmul模板参数如下：
84+ 
84```85```
85template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE = C_TYPE, const auto& MM_CFG = CFG_NORM, class MM_CB = MatmulCallBackFunc<nullptr, nullptr, nullptr>, MATMUL_POLICY_DEFAULT_OF(MatmulPolicy)>86template <class A_TYPE, class B_TYPE, class C_TYPE, class BIAS_TYPE = C_TYPE, const auto& MM_CFG = CFG_NORM, class MM_CB = MatmulCallBackFunc<nullptr, nullptr, nullptr>, MATMUL_POLICY_DEFAULT_OF(MatmulPolicy)>
86```87```
Mdocs/zh/api/SIMD-API/adv_api/cube_compute/Matmul_Kernel/Matmul_usage.md+1-0
@@ -125,6 +125,7 @@ Matmul的计算公式为：C = A \* B + Bias，其示意图如下。
125 125 
126 <!-- npu="310p" id86 -->126 <!-- npu="310p" id86 -->
127 Atlas推理系列产品AI Core上需要额外调用[SetLocalWorkspace](SetLocalWorkspace.md)接口设置计算所需的UB空间。127 Atlas推理系列产品AI Core上需要额外调用[SetLocalWorkspace](SetLocalWorkspace.md)接口设置计算所需的UB空间。
128+ 
128 ```129 ```
129 mm.SetLocalWorkspace(usedUbBufLen);130 mm.SetLocalWorkspace(usedUbBufLen);
130 ```131 ```
Mdocs/zh/api/SIMD-API/basic_api/cube_compute_ISASI/cube_compute_load/DataCopy_L1ToBiasTable.md+2-0
@@ -52,12 +52,14 @@ DataCopy数据搬运支持将矩阵计算用到的Bias参数从L1 Buffer移动
52 template <typename T>52 template <typename T>
53 __aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const uint32_t count)53 __aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const uint32_t count)
54 ```54 ```
55+ 
55- 连续或非连续搬运场景（高维切分）。56- 连续或非连续搬运场景（高维切分）。
56 57 
57 ```cpp58 ```cpp
58 template <typename T>59 template <typename T>
59 __aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const DataCopyParams& repeatParams)60 __aicore__ inline void DataCopy(const LocalTensor<T>& dst, const LocalTensor<T>& src, const DataCopyParams& repeatParams)
60 ```61 ```
62+ 
61- 连续或非连续搬运场景，支持源操作数和目的操作数数据类型不一致。63- 连续或非连续搬运场景，支持源操作数和目的操作数数据类型不一致。
62 64 
63 ```cpp65 ```cpp
Mdocs/zh/api/SIMD-API/basic_api/cube_compute_ISASI/cube_store_key_features/accompanying_quantization.md+1-0
@@ -19,6 +19,7 @@ VQF322B8_PRE, // Vector_Quant_Float32_2_B8：float量化成int8_t/uint8
19REQ8, // ReQuant_int8：int32_t重量化成int8_t/uint8_t，scalar量化19REQ8, // ReQuant_int8：int32_t重量化成int8_t/uint8_t，scalar量化
20VREQ8, // Vector_ReQuant_int8：int32_t重量化成int8_t/uint8_t，tensor量化20VREQ8, // Vector_ReQuant_int8：int32_t重量化成int8_t/uint8_t，tensor量化
21```21```
22+ 
22<!-- npu="950" id1 -->23<!-- npu="950" id1 -->
23除上述量化模式外，Ascend 950PR&950DT系列产品产品还额外支持以下量化模式：24除上述量化模式外，Ascend 950PR&950DT系列产品产品还额外支持以下量化模式：
24 25 
Mdocs/zh/api/SIMD-API/basic_api/debug_interface/onboard_print/DumpAccChkPoint.md+2-0
@@ -42,6 +42,7 @@
42```cpp42```cpp
43AscendC::DumpAccChkPoint(srcLocal, 5, 32, dataLen);43AscendC::DumpAccChkPoint(srcLocal, 5, 32, dataLen);
44```44```
45+ 
45> [!CAUTION]注意46> [!CAUTION]注意
46> 该接口主要用于调试分析，开启后会对算子性能产生一定影响，通常在调试阶段使用，生产环境建议关闭。<br>47> 该接口主要用于调试分析，开启后会对算子性能产生一定影响，通常在调试阶段使用，生产环境建议关闭。<br>
47> 默认情况下，调用该接口就会打印相关内容，开发者可以参考[关闭ASCENDC_DUMP说明](../disable_ascendc_dump.md)，按需关闭该接口功能。48> 默认情况下，调用该接口就会打印相关内容，开发者可以参考[关闭ASCENDC_DUMP说明](../disable_ascendc_dump.md)，按需关闭该接口功能。
@@ -114,6 +115,7 @@ AscendC::DumpAccChkPoint(srcLocal, index, countOff, dumpSize);
114```115```
115 116 
116打印结果如下：117打印结果如下：
118+ 
117```plain119```plain
118DumpTensor: desc=56, addr=0x40, data_type=float16, position=UB, dump_size=15120DumpTensor: desc=56, addr=0x40, data_type=float16, position=UB, dump_size=15
119[4.710938, 4.707031, 4.773438, 2.271484, 4.347656, 2.359375, 1.284180, 1.073242, 1.242188, 2.298828, 0.521973, 1.099609, 1.880859, 1.226562, 3.916016]121[4.710938, 4.707031, 4.773438, 2.271484, 4.347656, 2.359375, 1.284180, 1.073242, 1.242188, 2.298828, 0.521973, 1.099609, 1.880859, 1.226562, 3.916016]
Mdocs/zh/api/SIMD-API/basic_api/debug_interface/onboard_print/DumpTensor.md+2-0
@@ -41,11 +41,13 @@
41```cpp41```cpp
42AscendC::DumpTensor(srcLocal, 5, dataLen);42AscendC::DumpTensor(srcLocal, 5, dataLen);
43```43```
44+ 
44> [!CAUTION]注意45> [!CAUTION]注意
45> 该接口主要用于调试分析，开启后会对算子性能产生一定影响，通常在调试阶段使用，生产环境建议关闭。<br>46> 该接口主要用于调试分析，开启后会对算子性能产生一定影响，通常在调试阶段使用，生产环境建议关闭。<br>
46> 默认情况下，调用该接口就会打印相关内容，开发者可以参考[关闭ASCENDC_DUMP说明](../disable_ascendc_dump.md)，按需关闭该接口功能。47> 默认情况下，调用该接口就会打印相关内容，开发者可以参考[关闭ASCENDC_DUMP说明](../disable_ascendc_dump.md)，按需关闭该接口功能。
47 48 
48打印示例如下： 49打印示例如下：
50+ 
49```plain51```plain
50DumpTensor: desc=5, addr=0, data_type=float16, position=UB, dump_size=3252DumpTensor: desc=5, addr=0, data_type=float16, position=UB, dump_size=32
51[19.000000, 4.000000, 38.000000, 50.000000, 39.000000, 67.000000, 84.000000, 98.000000, 21.000000, 36.000000, 18.000000, 46.000000, 10.000000, 92.000000, 26.000000, 38.000000, 39.000000, 9.000000, 82.000000, 37.000000, 35.000000, 65.000000, 97.000000, 59.000000, 89.000000, 63.000000, 70.000000, 57.000000, 35.000000, 3.000000, 16.000000,53[19.000000, 4.000000, 38.000000, 50.000000, 39.000000, 67.000000, 84.000000, 98.000000, 21.000000, 36.000000, 18.000000, 46.000000, 10.000000, 92.000000, 26.000000, 38.000000, 39.000000, 9.000000, 82.000000, 37.000000, 35.000000, 65.000000, 97.000000, 59.000000, 89.000000, 63.000000, 70.000000, 57.000000, 35.000000, 3.000000, 16.000000,
Mdocs/zh/api/SIMD-API/basic_api/memory_vector_compute/data_padding/Duplicate.md+2-0
@@ -46,6 +46,7 @@ def Duplicate(scalarValue, dst, count):
46 for i in range(count): 46 for i in range(count):
47 dst[i] = scalarValue47 dst[i] = scalarValue
48```48```
49+ 
49<!-- npu="950" id26 -->50<!-- npu="950" id26 -->
50针对Ascend 950PR&950DT系列产品，为方便开发者使用，tensor前n个数据计算接口同时也支持直接传入LocalTensor，此时会将LocalTensor的第一个元素复制多次并填充到向量中。51针对Ascend 950PR&950DT系列产品，为方便开发者使用，tensor前n个数据计算接口同时也支持直接传入LocalTensor，此时会将LocalTensor的第一个元素复制多次并填充到向量中。
51<!-- end id26 -->52<!-- end id26 -->
@@ -193,6 +194,7 @@ def Duplicate(scalarValue, dst, count):
193<!-- npu="950" id25 -->194<!-- npu="950" id25 -->
194 195 
195- tensor前n个数据计算样例，源操作数为Tensor196- tensor前n个数据计算样例，源操作数为Tensor
197+ 
196 ```cpp198 ```cpp
197 AscendC::Duplicate<half>(dstLocal, srcLocal, srcDataSize);199 AscendC::Duplicate<half>(dstLocal, srcLocal, srcDataSize);
198 ```200 ```
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/aux_data_types/RegLayout.md+2-0
@@ -1,6 +1,7 @@
1# RegLayout<a name="ZH-CN_TOPIC_0000001956862301"></a>1# RegLayout<a name="ZH-CN_TOPIC_0000001956862301"></a>
2 2 
3源操作数和目的操作数位宽不同时，单条指令计算量以位宽更大的数据类型为准，RegLayout用于控制位宽小的元素在寄存器中的排布方式。更多特性说明请参考[Cast](../type_conversion/Cast.md)。3源操作数和目的操作数位宽不同时，单条指令计算量以位宽更大的数据类型为准，RegLayout用于控制位宽小的元素在寄存器中的排布方式。更多特性说明请参考[Cast](../type_conversion/Cast.md)。
4+ 
4```cpp5```cpp
5enum class RegLayout { 6enum class RegLayout {
6 UNKNOWN = -1, 7 UNKNOWN = -1,
@@ -46,6 +47,7 @@ enum class RegLayout {
46 ![](../../../../figures/src_dst_bitwidth_4to1.png "源操作数与目的操作数位宽比为4:1")47 ![](../../../../figures/src_dst_bitwidth_4to1.png "源操作数与目的操作数位宽比为4:1")
47 48 
48Cast在不同RegLayout取值下的数据排布伪代码如下，其中layout参数与枚举值的映射关系为：0对应RegLayout::ZERO，1对应RegLayout::ONE，2对应RegLayout::TWO，3对应RegLayout::THREE：49Cast在不同RegLayout取值下的数据排布伪代码如下，其中layout参数与枚举值的映射关系为：0对应RegLayout::ZERO，1对应RegLayout::ONE，2对应RegLayout::TWO，3对应RegLayout::THREE：
50+ 
49```python51```python
50import numpy as np52import numpy as np
51 53 
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/reg_data_load/LoadAlign_MaskReg.md+3-0
@@ -117,6 +117,7 @@ LoadAlign能够实现数据从Unified Buffer（UB）搬运至[MaskReg](../regist
117## 调用示例<a name="section642mcpsimp"></a>117## 调用示例<a name="section642mcpsimp"></a>
118 118 
119- 普通搬运接口119- 普通搬运接口
120+ 
120 ```cpp121 ```cpp
121 template <typename T>122 template <typename T>
122 __simd_vf__ inline void ComputeMode01(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)123 __simd_vf__ inline void ComputeMode01(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)
@@ -130,6 +131,7 @@ LoadAlign能够实现数据从Unified Buffer（UB）搬运至[MaskReg](../regist
130 ```131 ```
131 132 
132- PostUpdate扩展搬运接口133- PostUpdate扩展搬运接口
134+ 
133 ```cpp135 ```cpp
134 template <typename T>136 template <typename T>
135 __simd_vf__ inline void StoreAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)137 __simd_vf__ inline void StoreAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)
@@ -143,6 +145,7 @@ LoadAlign能够实现数据从Unified Buffer（UB）搬运至[MaskReg](../regist
143 ```145 ```
144 146 
145- 使用AddrReg存储偏移量接口147- 使用AddrReg存储偏移量接口
148+ 
146 ```cpp149 ```cpp
147 template <typename T>150 template <typename T>
148 __simd_vf__ inline void ComputeMode03(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)151 __simd_vf__ inline void ComputeMode03(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/reg_data_store/StoreAlign_MaskReg.md+3-0
@@ -103,6 +103,7 @@ StoreAlign能够实现数据从[MaskReg](../register_data_types/MaskReg.md)搬
103## 调用示例<a name="section642mcpsimp"></a>103## 调用示例<a name="section642mcpsimp"></a>
104 104 
105- 普通搬运接口105- 普通搬运接口
106+ 
106 ```cpp107 ```cpp
107 template <typename T>108 template <typename T>
108 __simd_vf__ inline void ComputeMode01(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)109 __simd_vf__ inline void ComputeMode01(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)
@@ -116,6 +117,7 @@ StoreAlign能够实现数据从[MaskReg](../register_data_types/MaskReg.md)搬
116 ```117 ```
117 118 
118- PostUpdate扩展搬运接口119- PostUpdate扩展搬运接口
120+ 
119 ```cpp121 ```cpp
120 template <typename T>122 template <typename T>
121 __simd_vf__ inline void StoreAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)123 __simd_vf__ inline void StoreAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)
@@ -129,6 +131,7 @@ StoreAlign能够实现数据从[MaskReg](../register_data_types/MaskReg.md)搬
129 ```131 ```
130 132 
131- 使用AddrReg存储偏移量接口133- 使用AddrReg存储偏移量接口
134+ 
132 ```cpp135 ```cpp
133 template <typename T>136 template <typename T>
134 __simd_vf__ inline void ComputeMode03(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)137 __simd_vf__ inline void ComputeMode03(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/reg_data_store/StoreUnAlign_continuous.md+3-0
@@ -197,6 +197,7 @@ __simd_vf__ inline void LoadUnAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr,
197## 调用示例<a name="section642mcpsimp"></a>197## 调用示例<a name="section642mcpsimp"></a>
198 198 
199- PostUpdate扩展搬运接口199- PostUpdate扩展搬运接口
200+ 
200 ```cpp201 ```cpp
201 template <typename T>202 template <typename T>
202 __simd_vf__ inline void StoreUnAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t postUpdateStride, uint16_t repeatTimes)203 __simd_vf__ inline void StoreUnAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t postUpdateStride, uint16_t repeatTimes)
@@ -214,6 +215,7 @@ __simd_vf__ inline void LoadUnAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr,
214 ```215 ```
215 216 
216- 使用AddrReg寄存器存储偏移量接口217- 使用AddrReg寄存器存储偏移量接口
218+ 
217 ```cpp219 ```cpp
218 template <typename T>220 template <typename T>
219 __simd_vf__ inline void StoreUnAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)221 __simd_vf__ inline void StoreUnAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)
@@ -233,6 +235,7 @@ __simd_vf__ inline void LoadUnAlignVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr,
233 ```235 ```
234 236 
235- 使用AR寄存器存储偏移量接口237- 使用AR寄存器存储偏移量接口
238+ 
236 ```cpp239 ```cpp
237 template <typename T>240 template <typename T>
238 __aicore__ inline void SqueezeVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)241 __aicore__ inline void SqueezeVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr, uint16_t oneRepeatSize, uint16_t repeatTimes)
Mdocs/zh/api/SIMD-API/basic_api/reg_vector_compute/register_data_types/MaskReg.md+4-0
@@ -38,11 +38,14 @@ MaskReg寄存器用于指示在计算过程中哪些元素参与计算，宽度
38## 函数原型<a name="section620mcpsimp"></a>38## 函数原型<a name="section620mcpsimp"></a>
39 39 
40- CreateMask接口40- CreateMask接口
41+ 
41 ```cpp42 ```cpp
42 template <typename T, MaskPattern mode = MaskPattern::ALL, const RegTrait& regTrait = RegTraitNumOne>43 template <typename T, MaskPattern mode = MaskPattern::ALL, const RegTrait& regTrait = RegTraitNumOne>
43 __simd_callee__ inline MaskReg CreateMask()44 __simd_callee__ inline MaskReg CreateMask()
44 ```45 ```
46+ 
45- UpdateMask接口47- UpdateMask接口
48+ 
46 ```cpp49 ```cpp
47 template <typename T, const RegTrait& regTrait = RegTraitNumOne>50 template <typename T, const RegTrait& regTrait = RegTraitNumOne>
48 __simd_callee__ inline MaskReg UpdateMask(uint32_t& scalarValue)51 __simd_callee__ inline MaskReg UpdateMask(uint32_t& scalarValue)
@@ -179,6 +182,7 @@ __simd_callee__ inline void Add(U& dstReg, U& srcReg0, U& srcReg1, MaskReg& mask
179 }182 }
180 }183 }
181 ```184 ```
185+ 
182 >[!NOTE]说明186 >[!NOTE]说明
183 >187 >
184 >当数据类型为int16_t（操作数为16位）时，每次循环内能够处理的元素个数Mask∈[1, 128]。当count>128时，仍按一次repeat（128元素）执行，不支持超过128元素的Mask控制。188 >当数据类型为int16_t（操作数为16位）时，每次循环内能够处理的元素个数Mask∈[1, 128]。当count>128时，仍按一次repeat（128元素）执行，不支持超过128元素的Mask控制。
Mdocs/zh/api/SIMD-API/basic_api/resource_management/TBufPool/InitBuffer.md+3-0
@@ -40,11 +40,14 @@
40## 函数原型<a name="section620mcpsimp"></a>40## 函数原型<a name="section620mcpsimp"></a>
41 41 
42- 为TQue分配内存42- 为TQue分配内存
43+ 
43 ```cpp44 ```cpp
44 template <class T>45 template <class T>
45 __aicore__ inline bool InitBuffer(T& que, uint8_t num, uint32_t len)46 __aicore__ inline bool InitBuffer(T& que, uint8_t num, uint32_t len)
46 ```47 ```
48+ 
47- 为TBuf分配内存49- 为TBuf分配内存
50+ 
48 ```cpp51 ```cpp
49 template <TPosition pos>52 template <TPosition pos>
50 __aicore__ inline bool InitBuffer(TBuf<pos>& buf, uint32_t len)53 __aicore__ inline bool InitBuffer(TBuf<pos>& buf, uint32_t len)
Mdocs/zh/api/SIMD-API/c_api/cache_ctrl/asc_set_scalar_cache_mode.md+0-1
@@ -82,7 +82,6 @@ PIPE_S
82```bash82```bash
83bisheng example.asc -o main --npu-arch=dav-3510 && ./main83bisheng example.asc -o main --npu-arch=dav-3510 && ./main
84```84```
85- 
86<!-- end id8 -->85<!-- end id8 -->
87 86 
88- 场景A：标量load使用`NORMAL_FIRST_VICTIM`读取GM tiling表。<a id="asc_set_scalar_cache_mode_scenario_a"></a>87- 场景A：标量load使用`NORMAL_FIRST_VICTIM`读取GM tiling表。<a id="asc_set_scalar_cache_mode_scenario_a"></a>
Mdocs/zh/api/SIMD-API/c_api/cube_compute/mmad_compute_key_features/GEMV.md+1-0
@@ -62,6 +62,7 @@ asc_mmad(c, a, b, 1, k, n, asc_unit_flag_mode::DISABLE,
62 62 
63<!-- npu="950" id3 -->63<!-- npu="950" id3 -->
64针对Ascend 950PR&950DT系列产品产品，可以通过设置asc_mmad的disable_gemv参数为true，将该功能关闭，示例如下：64针对Ascend 950PR&950DT系列产品产品，可以通过设置asc_mmad的disable_gemv参数为true，将该功能关闭，示例如下：
65+ 
65```cpp66```cpp
66asc_mmad(c, a, b, 1, k, n, asc_unit_flag_mode::DISABLE,67asc_mmad(c, a, b, 1, k, n, asc_unit_flag_mode::DISABLE,
67 true /* disable_gemv */, false, true);68 true /* disable_gemv */, false, true);
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_copy_l12bt/asc_copy_l12bt_arch_2201.md+2-0
@@ -33,11 +33,13 @@
33## 函数原型33## 函数原型
34 34 
35- 前n个数据搬运35- 前n个数据搬运
36+ 
36 ```cpp37 ```cpp
37 __aicore__ inline void asc_copy_l12bt(uint64_t dst, __cbuf__ void* src, uint32_t size)38 __aicore__ inline void asc_copy_l12bt(uint64_t dst, __cbuf__ void* src, uint32_t size)
38 ```39 ```
39 40 
40- 高维切分搬运41- 高维切分搬运
42+ 
41 ```cpp43 ```cpp
42 __aicore__ inline void asc_copy_l12bt(uint64_t dst, __cbuf__ void* src, uint16_t conv_control, uint16_t n_burst, uint16_t len_burst, uint16_t source_gap, uint16_t dst_gap)44 __aicore__ inline void asc_copy_l12bt(uint64_t dst, __cbuf__ void* src, uint16_t conv_control, uint16_t n_burst, uint16_t len_burst, uint16_t source_gap, uint16_t dst_gap)
43 ```45 ```
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_copy_l12fb/asc_copy_l12fb_arch_2201.md+2-0
@@ -33,11 +33,13 @@
33## 函数原型33## 函数原型
34 34 
35- 前n个数据搬运35- 前n个数据搬运
36+ 
36 ```cpp37 ```cpp
37 __aicore__ inline void asc_copy_l12fb(__fbuf__ void* dst, __cbuf__ void* src, uint32_t size)38 __aicore__ inline void asc_copy_l12fb(__fbuf__ void* dst, __cbuf__ void* src, uint32_t size)
38 ```39 ```
39 40 
40- 高维切分搬运41- 高维切分搬运
42+ 
41 ```cpp43 ```cpp
42 __aicore__ inline void asc_copy_l12fb(__fbuf__ void* dst, __cbuf__ void* src, uint16_t n_burst, uint16_t len_burst, uint16_t src_gap_size, uint16_t dst_gap_size)44 __aicore__ inline void asc_copy_l12fb(__fbuf__ void* dst, __cbuf__ void* src, uint16_t n_burst, uint16_t len_burst, uint16_t src_gap_size, uint16_t dst_gap_size)
43 ```45 ```
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_copy_l12l0a/asc_copy_l12l0a_arch_2201.md+1-0
@@ -33,6 +33,7 @@
33## 函数原型33## 函数原型
34 34 
35- 高维切分搬运2D格式35- 高维切分搬运2D格式
36+ 
36 ```cpp37 ```cpp
37 __aicore__ inline void asc_copy_l12l0a(__ca__ int4b_t* dst, __cbuf__ int4b_t* src, uint16_t start_index, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap)38 __aicore__ inline void asc_copy_l12l0a(__ca__ int4b_t* dst, __cbuf__ int4b_t* src, uint16_t start_index, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap)
38 __aicore__ inline void asc_copy_l12l0a(__ca__ uint8_t* dst, __cbuf__ uint8_t* src, uint16_t start_index, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap)39 __aicore__ inline void asc_copy_l12l0a(__ca__ uint8_t* dst, __cbuf__ uint8_t* src, uint16_t start_index, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap)
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_copy_l12l0b/asc_copy_l12l0b_arch_2201.md+1-0
@@ -33,6 +33,7 @@
33## 函数原型33## 函数原型
34 34 
35- 高维切分搬运2D格式35- 高维切分搬运2D格式
36+ 
36 ```cpp37 ```cpp
37 __aicore__ inline void asc_copy_l12l0b(__cb__ int4b_t* dst, __cbuf__ int4b_t* src, uint16_t start_index, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap)38 __aicore__ inline void asc_copy_l12l0b(__cb__ int4b_t* dst, __cbuf__ int4b_t* src, uint16_t start_index, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap)
38 __aicore__ inline void asc_copy_l12l0b(__cb__ uint8_t* dst, __cbuf__ uint8_t* src, uint16_t start_index, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap)39 __aicore__ inline void asc_copy_l12l0b(__cb__ uint8_t* dst, __cbuf__ uint8_t* src, uint16_t start_index, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap)
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_copy_l12l0b_sparse.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 高维切分搬运37- 高维切分搬运
38+ 
38 ```cpp39 ```cpp
39 __aicore__ inline void asc_copy_l12l0b_sparse(__cb__ int8_t* dst, __cbuf__ int8_t* src, __cbuf__ int8_t* index, uint16_t start_index, uint8_t repeat)40 __aicore__ inline void asc_copy_l12l0b_sparse(__cb__ int8_t* dst, __cbuf__ int8_t* src, __cbuf__ int8_t* index, uint16_t start_index, uint8_t repeat)
40 ```41 ```
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/asc_copy_l12l0b_trans/asc_copy_l12l0b_trans_arch_2201.md+1-0
@@ -78,6 +78,7 @@
78## 函数原型78## 函数原型
79 79 
80- 高维切分搬运80- 高维切分搬运
81+ 
81```cpp82```cpp
82__aicore__ inline void asc_copy_l12l0b_trans(__cb__ uint8_t* dst, __cbuf__ uint8_t* src, uint16_t index_id, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap, bool enable_addr_decrement, uint16_t dst_frac_gap)83__aicore__ inline void asc_copy_l12l0b_trans(__cb__ uint8_t* dst, __cbuf__ uint8_t* src, uint16_t index_id, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap, bool enable_addr_decrement, uint16_t dst_frac_gap)
83__aicore__ inline void asc_copy_l12l0b_trans(__cb__ int8_t* dst, __cbuf__ int8_t* src, uint16_t index_id, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap, bool enable_addr_decrement, uint16_t dst_frac_gap)84__aicore__ inline void asc_copy_l12l0b_trans(__cb__ int8_t* dst, __cbuf__ int8_t* src, uint16_t index_id, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap, bool enable_addr_decrement, uint16_t dst_frac_gap)
Mdocs/zh/api/SIMD-API/c_api/cube_datamove/cube_store_key_features/accompanying_quantization.md+1-0
@@ -41,6 +41,7 @@ VQF322B8_PRE, // Vector_Quant_Float32_2_B8：float量化成int8_t/uint8
41REQ8, // ReQuant_int8：int32_t重量化成int8_t/uint8_t，scalar量化41REQ8, // ReQuant_int8：int32_t重量化成int8_t/uint8_t，scalar量化
42VREQ8, // Vector_ReQuant_int8：int32_t重量化成int8_t/uint8_t，tensor量化42VREQ8, // Vector_ReQuant_int8：int32_t重量化成int8_t/uint8_t，tensor量化
43```43```
44+ 
44<!-- npu="950" id3 -->45<!-- npu="950" id3 -->
45除上述量化模式外，Ascend 950PR&950DT系列产品产品还额外支持以下量化模式：46除上述量化模式外，Ascend 950PR&950DT系列产品产品还额外支持以下量化模式：
46 47 
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_bfloat162float_sync_deprecated.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 同步计算37- 同步计算
38+ 
38```cpp39```cpp
39__aicore__ inline void asc_bfloat162float_sync(__ubuf__ float* dst, __ubuf__ bfloat16_t* src, uint32_t count)40__aicore__ inline void asc_bfloat162float_sync(__ubuf__ float* dst, __ubuf__ bfloat16_t* src, uint32_t count)
40```41```
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_bfloat162int32_sync_deprecated.md+1-0
@@ -41,6 +41,7 @@
41## 函数原型41## 函数原型
42 42 
43- 同步计算43- 同步计算
44+ 
44 ```cpp45 ```cpp
45 // RINT舍入模式46 // RINT舍入模式
46 __aicore__ inline void asc_bfloat162int32_rn_sync(__ubuf__ int32_t* dst, __ubuf__ bfloat16_t* src, uint32_t count)47 __aicore__ inline void asc_bfloat162int32_rn_sync(__ubuf__ int32_t* dst, __ubuf__ bfloat16_t* src, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_bitsort_sync_deprecated.md+1-0
@@ -37,6 +37,7 @@ Score和Index分别存储在src0和src1中，按Score进行排序（Score大的
37## 函数原型37## 函数原型
38 38 
39- 同步计算39- 同步计算
40+ 
40 ```cpp41 ```cpp
41 __aicore__ inline void asc_bitsort_sync(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ uint32_t* src1, int32_t repeat)42 __aicore__ inline void asc_bitsort_sync(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ uint32_t* src1, int32_t repeat)
42 __aicore__ inline void asc_bitsort_sync(__ubuf__ float* dst, __ubuf__ float* src0, __ubuf__ uint32_t* src1, int32_t repeat)43 __aicore__ inline void asc_bitsort_sync(__ubuf__ float* dst, __ubuf__ float* src0, __ubuf__ uint32_t* src1, int32_t repeat)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_copy_l12bt_sync_arch_2201_deprecated.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 同步搬运37- 同步搬运
38+ 
38 ```cpp39 ```cpp
39 __aicore__ inline void asc_copy_l12bt_sync(uint64_t dst, __cbuf__ void* src, uint32_t size)40 __aicore__ inline void asc_copy_l12bt_sync(uint64_t dst, __cbuf__ void* src, uint32_t size)
40 ```41 ```
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_copy_l12fb_sync_arch_2201_deprecated.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 同步搬运37- 同步搬运
38+ 
38 ```cpp39 ```cpp
39 __aicore__ inline void asc_copy_l12fb_sync(__fbuf__ void* dst, __cbuf__ void* src, uint32_t size)40 __aicore__ inline void asc_copy_l12fb_sync(__fbuf__ void* dst, __cbuf__ void* src, uint32_t size)
40 ```41 ```
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_copy_l12fb_sync_arch_3510_deprecated.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 同步搬运37- 同步搬运
38+ 
38 ```cpp39 ```cpp
39 __aicore__ inline void asc_copy_l12fb_sync(__fbuf__ void* dst, __cbuf__ void* src, uint32_t size)40 __aicore__ inline void asc_copy_l12fb_sync(__fbuf__ void* dst, __cbuf__ void* src, uint32_t size)
40 ```41 ```
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_copy_l12l0b_sparse_sync_deprecated.md+1-0
@@ -37,6 +37,7 @@
37## 函数原型37## 函数原型
38 38 
39- 同步搬运39- 同步搬运
40+ 
40 ```cpp41 ```cpp
41 __aicore__ inline void asc_copy_l12l0b_sparse_sync(__cb__ int8_t* dst, __cbuf__ int8_t* src, __cbuf__ int8_t* index, uint16_t start_index, uint8_t repeat)42 __aicore__ inline void asc_copy_l12l0b_sparse_sync(__cb__ int8_t* dst, __cbuf__ int8_t* src, __cbuf__ int8_t* index, uint16_t start_index, uint8_t repeat)
42 ```43 ```
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_copy_l12l0b_trans_sync_arch_2201_deprecated.md+1-0
@@ -80,6 +80,7 @@
80## 函数原型80## 函数原型
81 81 
82- 同步高维切分搬运82- 同步高维切分搬运
83+ 
83```cpp84```cpp
84__aicore__ inline void asc_copy_l12l0b_trans_sync(__cb__ uint8_t* dst, __cbuf__ uint8_t* src, uint16_t index_id, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap, bool enable_addr_decrement, uint16_t dst_frac_gap)85__aicore__ inline void asc_copy_l12l0b_trans_sync(__cb__ uint8_t* dst, __cbuf__ uint8_t* src, uint16_t index_id, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap, bool enable_addr_decrement, uint16_t dst_frac_gap)
85__aicore__ inline void asc_copy_l12l0b_trans_sync(__cb__ int8_t* dst, __cbuf__ int8_t* src, uint16_t index_id, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap, bool enable_addr_decrement, uint16_t dst_frac_gap)86__aicore__ inline void asc_copy_l12l0b_trans_sync(__cb__ int8_t* dst, __cbuf__ int8_t* src, uint16_t index_id, uint8_t repeat, uint16_t src_stride, uint16_t dst_gap, bool enable_addr_decrement, uint16_t dst_frac_gap)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_copy_l12ub_sync_deprecated.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 同步搬运37- 同步搬运
38+ 
38 ```cpp39 ```cpp
39 __aicore__ inline void asc_copy_l12ub_sync(__ubuf__ void* dst_addr, __cbuf__ void* src_addr, int8_t sub_blockid, uint16_t burst_count, uint16_t burst_len, uint16_t src_gap, uint16_t dst_gap)40 __aicore__ inline void asc_copy_l12ub_sync(__ubuf__ void* dst_addr, __cbuf__ void* src_addr, int8_t sub_blockid, uint16_t burst_count, uint16_t burst_len, uint16_t src_gap, uint16_t dst_gap)
40 ```41 ```
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_eq_sync_deprecated.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 同步计算37- 同步计算
38+ 
38 ```cpp39 ```cpp
39 // 比较结果存入寄存器40 // 比较结果存入寄存器
40 __aicore__ inline void asc_eq_sync(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)41 __aicore__ inline void asc_eq_sync(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_exp_sync_deprecated.md+1-0
@@ -38,6 +38,7 @@ $$
38## 函数原型38## 函数原型
39 39 
40- 同步计算40- 同步计算
41+ 
41```cpp42```cpp
42__aicore__ inline void asc_exp_sync(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)43__aicore__ inline void asc_exp_sync(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)
43__aicore__ inline void asc_exp_sync(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)44__aicore__ inline void asc_exp_sync(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_fma_sync_deprecated.md+1-0
@@ -38,6 +38,7 @@ $$
38## 函数原型38## 函数原型
39 39 
40- 同步计算40- 同步计算
41+ 
41 ```cpp42 ```cpp
42 __aicore__ inline void asc_fma_sync(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)43 __aicore__ inline void asc_fma_sync(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)
43 __aicore__ inline void asc_fma_sync(__ubuf__ float* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)44 __aicore__ inline void asc_fma_sync(__ubuf__ float* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_ge_scalar_sync_deprecated.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 同步计算37- 同步计算
38+ 
38 ```cpp39 ```cpp
39 __aicore__ inline void asc_ge_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)40 __aicore__ inline void asc_ge_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
40 __aicore__ inline void asc_ge_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)41 __aicore__ inline void asc_ge_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_gt_scalar_sync_deprecated.md+1-0
@@ -35,6 +35,7 @@ src中的每个元素逐个与标量value比较大小，如果某个位置上的
35## 函数原型35## 函数原型
36 36 
37- 同步计算37- 同步计算
38+ 
38 ```cpp39 ```cpp
39 __aicore__ inline void asc_gt_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)40 __aicore__ inline void asc_gt_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
40 __aicore__ inline void asc_gt_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)41 __aicore__ inline void asc_gt_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_half2int16_sync_deprecated.md+1-0
@@ -40,6 +40,7 @@
40## 函数原型40## 函数原型
41 41 
42- 同步计算42- 同步计算
43+ 
43 ```cpp44 ```cpp
44 // RINT舍入模式45 // RINT舍入模式
45 __aicore__ inline void asc_half2int16_rn_sync(__ubuf__ int16_t* dst, __ubuf__ half* src, uint32_t count)46 __aicore__ inline void asc_half2int16_rn_sync(__ubuf__ int16_t* dst, __ubuf__ half* src, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_le_scalar_sync_deprecated.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 同步计算37- 同步计算
38+ 
38 ```cpp39 ```cpp
39 __aicore__ inline void asc_le_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)40 __aicore__ inline void asc_le_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
40 __aicore__ inline void asc_le_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)41 __aicore__ inline void asc_le_scalar_sync(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_le_sync_deprecated.md+1-0
@@ -36,6 +36,7 @@
36## 函数原型36## 函数原型
37 37 
38- 同步计算38- 同步计算
39+ 
39 ```cpp40 ```cpp
40 // 结果输出到寄存器中41 // 结果输出到寄存器中
41 __aicore__ inline void asc_le_sync(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)42 __aicore__ inline void asc_le_sync(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_log_sync_deprecated.md+1-0
@@ -38,6 +38,7 @@ $$
38## 函数原型38## 函数原型
39 39 
40- 同步计算40- 同步计算
41+ 
41 ```cpp42 ```cpp
42 __aicore__ inline void asc_log_sync(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)43 __aicore__ inline void asc_log_sync(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)
43 __aicore__ inline void asc_log_sync(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)44 __aicore__ inline void asc_log_sync(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_min_sync_deprecated.md+1-0
@@ -38,6 +38,7 @@ $$
38## 函数原型38## 函数原型
39 39 
40- 同步计算40- 同步计算
41+ 
41```cpp42```cpp
42__aicore__ inline void asc_min_sync(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint32_t count)43__aicore__ inline void asc_min_sync(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint32_t count)
43__aicore__ inline void asc_min_sync(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)44__aicore__ inline void asc_min_sync(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_ne_sync_deprecated.md+1-0
@@ -36,6 +36,7 @@
36## 函数原型36## 函数原型
37 37 
38- 同步计算38- 同步计算
39+ 
39 ```cpp40 ```cpp
40 // 结果输出到寄存器中41 // 结果输出到寄存器中
41 __aicore__ inline void asc_ne_sync(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)42 __aicore__ inline void asc_ne_sync(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_not_sync_deprecated.md+1-0
@@ -38,6 +38,7 @@ $$
38## 函数原型38## 函数原型
39 39 
40- 同步计算40- 同步计算
41+ 
41 ```cpp42 ```cpp
42 __aicore__ inline void asc_not_sync(__ubuf__ int16_t* dst, __ubuf__ int16_t* src, uint32_t count)43 __aicore__ inline void asc_not_sync(__ubuf__ int16_t* dst, __ubuf__ int16_t* src, uint32_t count)
43 __aicore__ inline void asc_not_sync(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src, uint32_t count)44 __aicore__ inline void asc_not_sync(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_or_sync_deprecated.md+1-0
@@ -38,6 +38,7 @@ $$
38## 函数原型38## 函数原型
39 39 
40- 同步计算40- 同步计算
41+ 
41 ```cpp42 ```cpp
42 __aicore__ inline void asc_or_sync(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint32_t count)43 __aicore__ inline void asc_or_sync(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint32_t count)
43 __aicore__ inline void asc_or_sync(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src0, __ubuf__ uint16_t* src1, uint32_t count)44 __aicore__ inline void asc_or_sync(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src0, __ubuf__ uint16_t* src1, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_relu_sync_deprecated.md+1-0
@@ -38,6 +38,7 @@ $$
38## 函数原型38## 函数原型
39 39 
40- 同步计算40- 同步计算
41+ 
41 ```cpp42 ```cpp
42 __aicore__ inline void asc_relu_sync(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)43 __aicore__ inline void asc_relu_sync(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)
43 __aicore__ inline void asc_relu_sync(__ubuf__ int32_t* dst, __ubuf__ int32_t* src, uint32_t count)44 __aicore__ inline void asc_relu_sync(__ubuf__ int32_t* dst, __ubuf__ int32_t* src, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_sqrt_sync_deprecated.md+1-0
@@ -38,6 +38,7 @@ $$
38## 函数原型38## 函数原型
39 39 
40- 同步计算40- 同步计算
41+ 
41 ```cpp42 ```cpp
42 __aicore__ inline void asc_sqrt_sync(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)43 __aicore__ inline void asc_sqrt_sync(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)
43 __aicore__ inline void asc_sqrt_sync(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)44 __aicore__ inline void asc_sqrt_sync(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_transpose_sync_deprecated.md+1-0
@@ -35,6 +35,7 @@
35## 函数原型35## 函数原型
36 36 
37- 同步计算37- 同步计算
38+ 
38 ```cpp39 ```cpp
39 __aicore__ inline void asc_transpose_sync(__ubuf__ int16_t* dst, __ubuf__ int16_t* src)40 __aicore__ inline void asc_transpose_sync(__ubuf__ int16_t* dst, __ubuf__ int16_t* src)
40 __aicore__ inline void asc_transpose_sync(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src)41 __aicore__ inline void asc_transpose_sync(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src)
Mdocs/zh/api/SIMD-API/c_api/deprecated_interface/asc_transto5hd_sync_deprecated.md+1-0File content is under review, please try again later
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign.md+1-0
@@ -220,6 +220,7 @@ __simd_callee__ inline void asc_loadalign(vector_float& dst,
220 220 
221- 该模式不支持目的操作数为掩码寄存器。221- 该模式不支持目的操作数为掩码寄存器。
222- 当dtype为`int64_t`时，由于掩码创建接口不支持b64模式，需要先通过[asc_create_mask_b32](../reg_mask/asc_create_mask.md)生成b32掩码，再通过[asc_unpack_lower](../reg_permute_sel/asc_unpack.md)将其展开为适用于b64数据的掩码后传入本接口。展开后的掩码以连续8个bit为一组，仅每组最低位的bit有效，用于控制对应的一个b64元素。222- 当dtype为`int64_t`时，由于掩码创建接口不支持b64模式，需要先通过[asc_create_mask_b32](../reg_mask/asc_create_mask.md)生成b32掩码，再通过[asc_unpack_lower](../reg_permute_sel/asc_unpack.md)将其展开为适用于b64数据的掩码后传入本接口。展开后的掩码以连续8个bit为一组，仅每组最低位的bit有效，用于控制对应的一个b64元素。
223+ 
223 ```c224 ```c
224 vector_int64_t src_reg;225 vector_int64_t src_reg;
225 vector_bool mask_b32 = asc_create_mask_b32(PAT_ALL);226 vector_bool mask_b32 = asc_create_mask_b32(PAT_ALL);
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadalign_postupdate.md+1-0
@@ -135,6 +135,7 @@ __simd_callee__ inline void asc_loadalign_postupdate(vector_float& dst,
135 135 
136- 该模式不支持目的操作数为掩码寄存器。136- 该模式不支持目的操作数为掩码寄存器。
137- 当dtype为`int64_t`时，由于掩码创建接口不支持b64模式，需要先通过[asc_create_mask_b32](../reg_mask/asc_create_mask.md)生成b32掩码，再通过[asc_unpack_lower](../reg_permute_sel/asc_unpack.md)将其展开为适用于b64数据的掩码后传入本接口。展开后的掩码以连续8个bit为一组，仅每组最低位的bit有效，用于控制对应的一个b64元素。137- 当dtype为`int64_t`时，由于掩码创建接口不支持b64模式，需要先通过[asc_create_mask_b32](../reg_mask/asc_create_mask.md)生成b32掩码，再通过[asc_unpack_lower](../reg_permute_sel/asc_unpack.md)将其展开为适用于b64数据的掩码后传入本接口。展开后的掩码以连续8个bit为一组，仅每组最低位的bit有效，用于控制对应的一个b64元素。
138+ 
138 ```c139 ```c
139 vector_int64_t src_reg;140 vector_int64_t src_reg;
140 vector_bool mask_b32 = asc_create_mask_b32(PAT_ALL);141 vector_bool mask_b32 = asc_create_mask_b32(PAT_ALL);
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadunalign_postupdate.md+1-0
@@ -138,6 +138,7 @@ __simd_callee__ inline void asc_loadunalign_postupdate(vector_int8_t& dst,
138- 每次循环迭代中，必须先调用`asc_update_addr_reg`生成该次迭代的起始偏移，再调用[asc_loadunalign_pre](asc_loadunalign_pre.md)初始化`src0`。该初始化调用与紧随其后的首次搬入必须使用相同的`src1`和`offset`。138- 每次循环迭代中，必须先调用`asc_update_addr_reg`生成该次迭代的起始偏移，再调用[asc_loadunalign_pre](asc_loadunalign_pre.md)初始化`src0`。该初始化调用与紧随其后的首次搬入必须使用相同的`src1`和`offset`。
139- 同一循环迭代内，后续搬入继续使用上一次本接口调用按`inc`更新后的`offset`时，可复用同步更新后的`src0`。如果`src1`或`offset`被另行修改，必须根据修改后的实际读取地址重新初始化`src0`。139- 同一循环迭代内，后续搬入继续使用上一次本接口调用按`inc`更新后的`offset`时，可复用同步更新后的`src0`。如果`src1`或`offset`被另行修改，必须根据修改后的实际读取地址重新初始化`src0`。
140- 当dtype为`int64_t`时，由于地址寄存器更新接口不支持b64模式，需要使用[asc_update_addr_reg_b32](../reg_addr_reg/asc_update_addr_reg.md)生成地址寄存器。一个b64元素占用两个b32寻址单元，因此传入`asc_update_addr_reg_b32`的各维偏移量以及本接口的`inc`都必须设置为期望的b64元素偏移量的2倍。140- 当dtype为`int64_t`时，由于地址寄存器更新接口不支持b64模式，需要使用[asc_update_addr_reg_b32](../reg_addr_reg/asc_update_addr_reg.md)生成地址寄存器。一个b64元素占用两个b32寻址单元，因此传入`asc_update_addr_reg_b32`的各维偏移量以及本接口的`inc`都必须设置为期望的b64元素偏移量的2倍。
141+ 
141 ```c142 ```c
142 vector_int64_t src_reg;143 vector_int64_t src_reg;
143 vector_bool mask_b32 = asc_create_mask_b32(PAT_ALL);144 vector_bool mask_b32 = asc_create_mask_b32(PAT_ALL);
Mdocs/zh/api/SIMD-API/c_api/reg_compute/load/asc_loadunalign_pre.md+1-0
@@ -129,6 +129,7 @@ __simd_callee__ inline void asc_loadunalign_pre(vector_load_unalign& dst,
129- 本接口和紧随其后的首次搬入必须使用相同的`src`和`offset`。129- 本接口和紧随其后的首次搬入必须使用相同的`src`和`offset`。
130- 同一循环迭代内，后续搬入继续使用[asc_loadunalign_postupdate](asc_loadunalign_postupdate.md)按`inc`更新后的`offset`时，可以复用更新后的`src0`；如果`src`或`offset`被另行修改，必须根据修改后的实际读取地址重新初始化`src0`。130- 同一循环迭代内，后续搬入继续使用[asc_loadunalign_postupdate](asc_loadunalign_postupdate.md)按`inc`更新后的`offset`时，可以复用更新后的`src0`；如果`src`或`offset`被另行修改，必须根据修改后的实际读取地址重新初始化`src0`。
131- 当dtype为`int64_t`时，由于地址寄存器更新接口不支持b64模式，需要使用[asc_update_addr_reg_b32](../reg_addr_reg/asc_update_addr_reg.md)生成地址寄存器。一个b64元素占用两个b32寻址单元，因此传入`asc_update_addr_reg_b32`的各维偏移量必须设置为期望的b64元素偏移量的2倍。131- 当dtype为`int64_t`时，由于地址寄存器更新接口不支持b64模式，需要使用[asc_update_addr_reg_b32](../reg_addr_reg/asc_update_addr_reg.md)生成地址寄存器。一个b64元素占用两个b32寻址单元，因此传入`asc_update_addr_reg_b32`的各维偏移量必须设置为期望的b64元素偏移量的2倍。
132+ 
132 ```c133 ```c
133 vector_int64_t src_reg;134 vector_int64_t src_reg;
134 vector_bool mask_b32 = asc_create_mask_b32(PAT_ALL);135 vector_bool mask_b32 = asc_create_mask_b32(PAT_ALL);
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_addr_reg/asc_update_addr_reg.md+0-1
@@ -125,7 +125,6 @@
125```bash125```bash
126bisheng example.asc -o main --npu-arch=dav-3510 && ./main126bisheng example.asc -o main --npu-arch=dav-3510 && ./main
127```127```
128- 
129<!-- end id8 -->128<!-- end id8 -->
130 129 
131```cpp130```cpp
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/asc_half2float.md+1-0
@@ -91,6 +91,7 @@ __simd_callee__ inline vector_float asc_half2float(vector_half src,
91bisheng example.asc -o main --npu-arch=dav-3510 && ./main91bisheng example.asc -o main --npu-arch=dav-3510 && ./main
92```92```
93<!-- end id8 -->93<!-- end id8 -->
94+ 
94```c95```c
95#include <cstdint>96#include <cstdint>
96#include <iostream>97#include <iostream>
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/asc_half2hif8.md+1-0
@@ -115,6 +115,7 @@ __simd_callee__ inline vector_hifloat8_t asc_half2hif8_rna(vector_half src,
115bisheng example.asc -o main --npu-arch=dav-3510 && ./main115bisheng example.asc -o main --npu-arch=dav-3510 && ./main
116```116```
117<!-- end id8 -->117<!-- end id8 -->
118+ 
118```c119```c
119#include <cstdint>120#include <cstdint>
120#include <iostream>121#include <iostream>
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/asc_half2int32.md+1-0
@@ -108,6 +108,7 @@ __simd_callee__ inline vector_int32_t asc_half2int32_rd(vector_half src,
108bisheng example.asc -o main --npu-arch=dav-3510 && ./main108bisheng example.asc -o main --npu-arch=dav-3510 && ./main
109```109```
110<!-- end id8 -->110<!-- end id8 -->
111+ 
111```c112```c
112#include <cstdint>113#include <cstdint>
113#include <iostream>114#include <iostream>
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/asc_half2int4x2.md+1-0
@@ -107,6 +107,7 @@ __simd_callee__ inline vector_int4x2_t asc_half2int4x2_<round_mode>_sat(
107bisheng example.asc -o main --npu-arch=dav-3510 && ./main107bisheng example.asc -o main --npu-arch=dav-3510 && ./main
108```108```
109<!-- end id8 -->109<!-- end id8 -->
110+ 
110```c111```c
111#include <cstdint>112#include <cstdint>
112#include <iostream>113#include <iostream>
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/asc_half2int8.md+1-0
@@ -115,6 +115,7 @@ __simd_callee__ inline vector_int8_t asc_half2int8_rn(vector_half src,
115bisheng example.asc -o main --npu-arch=dav-3510 && ./main115bisheng example.asc -o main --npu-arch=dav-3510 && ./main
116```116```
117<!-- end id8 -->117<!-- end id8 -->
118+ 
118```c119```c
119#include <cstdint>120#include <cstdint>
120#include <iostream>121#include <iostream>
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/asc_half2uint8.md+1-0
@@ -113,6 +113,7 @@ __simd_callee__ inline vector_uint8_t asc_half2uint8_rd(vector_half src,
113bisheng example.asc -o main --npu-arch=dav-3510 && ./main113bisheng example.asc -o main --npu-arch=dav-3510 && ./main
114```114```
115<!-- end id8 -->115<!-- end id8 -->
116+ 
116```c117```c
117#include <cstdint>118#include <cstdint>
118#include <iostream>119#include <iostream>
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/asc_int4x22bfloat16.md+1-0
@@ -89,6 +89,7 @@ __simd_callee__ inline vector_bfloat16_t asc_int4x22bfloat16(
89bisheng example.asc -o main --npu-arch=dav-3510 && ./main89bisheng example.asc -o main --npu-arch=dav-3510 && ./main
90```90```
91<!-- end id8 -->91<!-- end id8 -->
92+ 
92```c93```c
93#include <cstdint>94#include <cstdint>
94#include <iostream>95#include <iostream>
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/asc_int4x22half.md+1-0
@@ -89,6 +89,7 @@ __simd_callee__ inline vector_half asc_int4x22half(
89bisheng example.asc -o main --npu-arch=dav-3510 && ./main89bisheng example.asc -o main --npu-arch=dav-3510 && ./main
90```90```
91<!-- end id8 -->91<!-- end id8 -->
92+ 
92```c93```c
93#include <cstdint>94#include <cstdint>
94#include <iostream>95#include <iostream>
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/asc_int4x22int16.md+1-0
@@ -89,6 +89,7 @@ __simd_callee__ inline vector_int16_t asc_int4x22int16(
89bisheng example.asc -o main --npu-arch=dav-3510 && ./main89bisheng example.asc -o main --npu-arch=dav-3510 && ./main
90```90```
91<!-- end id8 -->91<!-- end id8 -->
92+ 
92```c93```c
93#include <cstdint>94#include <cstdint>
94#include <iostream>95#include <iostream>
Mdocs/zh/api/SIMD-API/c_api/reg_compute/reg_convert/rounding_mode.md+13-0
@@ -33,6 +33,7 @@ Ascend C API支持以下舍入模式：
33- 当小数部分等于0.5时，舍入到最近的偶数33- 当小数部分等于0.5时，舍入到最近的偶数
34 34 
35**示例：**35**示例：**
36+ 
36```37```
371.4 → 1 (小数部分 < 0.5，向下)381.4 → 1 (小数部分 < 0.5，向下)
381.5 → 2 (小数部分 = 0.5，舍入到偶数)391.5 → 2 (小数部分 = 0.5，舍入到偶数)
@@ -48,6 +49,7 @@ Ascend C API支持以下舍入模式：
48- 当小数部分大于等于0.5时，向上舍入49- 当小数部分大于等于0.5时，向上舍入
49 50 
50**示例：**51**示例：**
52+ 
51```53```
521.4 → 1 (小数部分 < 0.5)541.4 → 1 (小数部分 < 0.5)
531.5 → 2 (小数部分 ≥ 0.5)551.5 → 2 (小数部分 ≥ 0.5)
@@ -60,6 +62,7 @@ Ascend C API支持以下舍入模式：
60始终向负无穷方向舍入，即使结果变小。62始终向负无穷方向舍入，即使结果变小。
61 63 
62**示例：**64**示例：**
65+ 
63```66```
641.7 → 1 (向下)671.7 → 1 (向下)
65-1.3 → -2 (向负无穷方向，变得更小)68-1.3 → -2 (向负无穷方向，变得更小)
@@ -70,6 +73,7 @@ Ascend C API支持以下舍入模式：
70始终向正无穷方向舍入，即使结果变大。73始终向正无穷方向舍入，即使结果变大。
71 74 
72**示例：**75**示例：**
76+ 
73```77```
741.3 → 2 (向上)781.3 → 2 (向上)
75-1.7 → -1 (向正无穷方向，变得更大)79-1.7 → -1 (向正无穷方向，变得更大)
@@ -80,6 +84,7 @@ Ascend C API支持以下舍入模式：
80直接截断小数部分，向零方向舍入。84直接截断小数部分，向零方向舍入。
81 85 
82**示例：**86**示例：**
87+ 
83```88```
841.7 → 1 (向零方向)891.7 → 1 (向零方向)
85-1.7 → -1 (向零方向)90-1.7 → -1 (向零方向)
@@ -90,6 +95,7 @@ Ascend C API支持以下舍入模式：
90舍入到最近的奇数，主要用于某些特殊的数值计算场景。95舍入到最近的奇数，主要用于某些特殊的数值计算场景。
91 96 
92**示例：**97**示例：**
98+ 
93```99```
941.5 → 1 (最近的奇数)1001.5 → 1 (最近的奇数)
952.5 → 3 (最近的奇数)1012.5 → 3 (最近的奇数)
@@ -114,6 +120,7 @@ Ascend C API支持以下舍入模式：
114- **二进制截断示例**：将int32转换为int8时，只保留低8位120- **二进制截断示例**：将int32转换为int8时，只保留低8位
115 121 
116**二进制截断示例（int32 → int8）：**122**二进制截断示例（int32 → int8）：**
123+ 
117```124```
118输入值（int32）: 0x000001FF (511)125输入值（int32）: 0x000001FF (511)
119二进制表示: 0000 0000 0000 0000 0000 0001 1111 1111126二进制表示: 0000 0000 0000 0000 0000 0001 1111 1111
@@ -138,6 +145,7 @@ Ascend C API支持以下舍入模式：
138- **NaN输入**：返回NaN145- **NaN输入**：返回NaN
139 146 
140**示例（float → half）：**147**示例（float → half）：**
148+ 
141```149```
142输入值: 70000.0 (超出half范围)150输入值: 70000.0 (超出half范围)
143输出值: +inf151输出值: +inf
@@ -156,6 +164,7 @@ Ascend C API支持以下舍入模式：
156- **NaN输入**：返回0164- **NaN输入**：返回0
157 165 
158**示例（float → int16）：**166**示例（float → int16）：**
167+ 
159```168```
160输入值: 100000.5 (超出int16范围)169输入值: 100000.5 (超出int16范围)
161非饱和模式: 结果被截断（具体值取决于二进制表示）170非饱和模式: 结果被截断（具体值取决于二进制表示）
@@ -177,6 +186,7 @@ Ascend C API支持以下舍入模式：
177- **保持数值安全**：不会出现因截断导致的意外结果186- **保持数值安全**：不会出现因截断导致的意外结果
178 187 
179**示例（int32 → int8）：**188**示例（int32 → int8）：**
189+ 
180```190```
181输入值: 511 (超出int8范围)191输入值: 511 (超出int8范围)
182饱和输出: 127 (int8最大值)192饱和输出: 127 (int8最大值)
@@ -189,6 +199,7 @@ Ascend C API支持以下舍入模式：
189```199```
190 200 
191**示例（int32 → uint8）：**201**示例（int32 → uint8）：**
202+ 
192```203```
193输入值: 500 (超出uint8范围)204输入值: 500 (超出uint8范围)
194饱和输出: 255 (uint8最大值)205饱和输出: 255 (uint8最大值)
@@ -207,6 +218,7 @@ Ascend C API支持以下舍入模式：
207- **NaN输入**：返回0218- **NaN输入**：返回0
208 219 
209**示例（float → half）：**220**示例（float → half）：**
221+ 
210```222```
211输入值: 70000.0 (超出half范围)223输入值: 70000.0 (超出half范围)
212饱和输出: 65504.0 (half最大值)224饱和输出: 65504.0 (half最大值)
@@ -228,6 +240,7 @@ Ascend C API支持以下舍入模式：
228- **NaN输入**：返回0240- **NaN输入**：返回0
229 241 
230**示例（float → int16）：**242**示例（float → int16）：**
243+ 
231```244```
232输入值: 100000.5 (超出int16范围)245输入值: 100000.5 (超出int16范围)
233饱和输出: 32767 (int16最大值)246饱和输出: 32767 (int16最大值)
Mdocs/zh/api/SIMD-API/c_api/sync/inter_core_sync/asc_sync_block_arrive.md+0-1
@@ -108,7 +108,6 @@ PIPE_S
108```bash108```bash
109bisheng example.asc -o main --npu-arch=dav-3510 && ./main109bisheng example.asc -o main --npu-arch=dav-3510 && ./main
110```110```
111- 
112<!-- end id8 -->111<!-- end id8 -->
113 112 
114```cpp113```cpp
Mdocs/zh/api/SIMD-API/c_api/sync/inter_core_sync/asc_sync_inter_arrive.md+0-1
@@ -108,7 +108,6 @@ PIPE_S
108```bash108```bash
109bisheng example.asc -o main --npu-arch=dav-3510 && ./main109bisheng example.asc -o main --npu-arch=dav-3510 && ./main
110```110```
111- 
112<!-- end id17 -->111<!-- end id17 -->
113 112 
114```cpp113```cpp
Mdocs/zh/api/SIMD-API/c_api/sync/inter_core_sync/asc_sync_intra_arrive.md+0-1
@@ -109,7 +109,6 @@ PIPE_S
109```bash109```bash
110bisheng example.asc -o main --npu-arch=dav-3510 && ./main110bisheng example.asc -o main --npu-arch=dav-3510 && ./main
111```111```
112- 
113<!-- end id8 -->112<!-- end id8 -->
114 113 
115```cpp114```cpp
Mdocs/zh/api/SIMD-API/c_api/sync/inter_core_sync/asc_sync_subblock_arrive.md+0-1
@@ -83,7 +83,6 @@ PIPE_S
83```bash83```bash
84bisheng example.asc -o main --npu-arch=dav-3510 && ./main84bisheng example.asc -o main --npu-arch=dav-3510 && ./main
85```85```
86- 
87<!-- end id17 -->86<!-- end id17 -->
88 87 
89```cpp88```cpp
Mdocs/zh/api/SIMD-API/c_api/sync/intra_core_sync/asc_lock.md+0-1
@@ -147,7 +147,6 @@ PIPE_S
147```bash147```bash
148bisheng example.asc -o main --npu-arch=dav-3510 && ./main148bisheng example.asc -o main --npu-arch=dav-3510 && ./main
149```149```
150- 
151<!-- end id8 -->150<!-- end id8 -->
152 151 
153```cpp152```cpp
Mdocs/zh/api/SIMD-API/c_api/sync/intra_core_sync/asc_sync.md+0-1
@@ -64,7 +64,6 @@ PIPE_S
64```bash64```bash
65bisheng example.asc -o main --npu-arch=dav-3510 && ./main65bisheng example.asc -o main --npu-arch=dav-3510 && ./main
66```66```
67- 
68<!-- end id8 -->67<!-- end id8 -->
69 68 
70```cpp69```cpp
Mdocs/zh/api/SIMD-API/c_api/sync/intra_core_sync/asc_sync_mte2.md+0-1
@@ -71,7 +71,6 @@ PIPE_S
71```bash71```bash
72bisheng example.asc -o main --npu-arch=dav-3510 && ./main72bisheng example.asc -o main --npu-arch=dav-3510 && ./main
73```73```
74- 
75<!-- end id8 -->74<!-- end id8 -->
76 75 
77```cpp76```cpp
Mdocs/zh/api/SIMD-API/c_api/sync/intra_core_sync/asc_sync_mte3.md+0-1
@@ -71,7 +71,6 @@ PIPE_S
71```bash71```bash
72bisheng example.asc -o main --npu-arch=dav-3510 && ./main72bisheng example.asc -o main --npu-arch=dav-3510 && ./main
73```73```
74- 
75<!-- end id8 -->74<!-- end id8 -->
76 75 
77```cpp76```cpp
Mdocs/zh/api/SIMD-API/c_api/sync/intra_core_sync/asc_sync_notify.md+0-1File content is under review, please try again later
Mdocs/zh/api/SIMD-API/c_api/sync/intra_core_sync/asc_sync_pipe.md+0-1File content is under review, please try again later
Mdocs/zh/api/SIMD-API/c_api/sync/intra_core_sync/asc_sync_vec.md+0-1
@@ -75,7 +75,6 @@ PIPE_S
75```bash75```bash
76bisheng example.asc -o main --npu-arch=dav-3510 && ./main76bisheng example.asc -o main --npu-arch=dav-3510 && ./main
77```77```
78- 
79<!-- end id8 -->78<!-- end id8 -->
80 79 
81```cpp80```cpp
Mdocs/zh/api/SIMD-API/c_api/utils/sys_misc/asc_nop.md+1-0
@@ -51,6 +51,7 @@ PIPE_S
51## 约束说明51## 约束说明
52 52 
53如果需要精确阻塞n拍流水，需要在n个`asc_nop`接口调用前后插入[`asc_sync`](../../sync/intra_core_sync/asc_sync.md)，例如：53如果需要精确阻塞n拍流水，需要在n个`asc_nop`接口调用前后插入[`asc_sync`](../../sync/intra_core_sync/asc_sync.md)，例如：
54+ 
54```c55```c
55// 阻塞3拍流水56// 阻塞3拍流水
56asc_sync();57asc_sync();
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_arith/asc_exp.md+2-0
@@ -36,12 +36,14 @@ $$
36## 函数原型36## 函数原型
37 37 
38- 前n个数据计算38- 前n个数据计算
39+ 
39```cpp40```cpp
40__aicore__ inline void asc_exp(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)41__aicore__ inline void asc_exp(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)
41__aicore__ inline void asc_exp(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)42__aicore__ inline void asc_exp(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)
42```43```
43 44 
44- 高维切分计算45- 高维切分计算
46+ 
45```cpp47```cpp
46__aicore__ inline void asc_exp(__ubuf__ half* dst, __ubuf__ half* src, uint8_t repeat,48__aicore__ inline void asc_exp(__ubuf__ half* dst, __ubuf__ half* src, uint8_t repeat,
47 uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)49 uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_arith/asc_log.md+2-0
@@ -36,12 +36,14 @@ $$
36## 函数原型36## 函数原型
37 37 
38- 前n个数据计算38- 前n个数据计算
39+ 
39 ```cpp40 ```cpp
40 __aicore__ inline void asc_log(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)41 __aicore__ inline void asc_log(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)
41 __aicore__ inline void asc_log(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)42 __aicore__ inline void asc_log(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)
42 ```43 ```
43 44 
44- 高维切分计算45- 高维切分计算
46+ 
45 ```cpp47 ```cpp
46 __aicore__ inline void asc_log(__ubuf__ half* dst, __ubuf__ half* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)48 __aicore__ inline void asc_log(__ubuf__ half* dst, __ubuf__ half* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
47 __aicore__ inline void asc_log(__ubuf__ float* dst, __ubuf__ float* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)49 __aicore__ inline void asc_log(__ubuf__ float* dst, __ubuf__ float* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_arith/asc_min.md+2-0
@@ -36,6 +36,7 @@ $$
36## 函数原型36## 函数原型
37 37 
38- 前n个数据计算38- 前n个数据计算
39+ 
39```cpp40```cpp
40__aicore__ inline void asc_min(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint32_t count)41__aicore__ inline void asc_min(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint32_t count)
41__aicore__ inline void asc_min(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)42__aicore__ inline void asc_min(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)
@@ -44,6 +45,7 @@ __aicore__ inline void asc_min(__ubuf__ float* dst, __ubuf__ float* src0, __ubuf
44```45```
45 46 
46- 高维切分计算47- 高维切分计算
48+ 
47```cpp49```cpp
48__aicore__ inline void asc_min(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint8_t repeat,50__aicore__ inline void asc_min(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint8_t repeat,
49 uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride,51 uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride,
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_arith/asc_sqrt.md+2-0
@@ -36,12 +36,14 @@ $$
36## 函数原型36## 函数原型
37 37 
38- 前n个数据计算38- 前n个数据计算
39+ 
39 ```cpp40 ```cpp
40 __aicore__ inline void asc_sqrt(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)41 __aicore__ inline void asc_sqrt(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)
41 __aicore__ inline void asc_sqrt(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)42 __aicore__ inline void asc_sqrt(__ubuf__ float* dst, __ubuf__ float* src, uint32_t count)
42 ```43 ```
43 44 
44- 高维切分计算45- 高维切分计算
46+ 
45 ```cpp47 ```cpp
46 __aicore__ inline void asc_sqrt(__ubuf__ half* dst, __ubuf__ half* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)48 __aicore__ inline void asc_sqrt(__ubuf__ half* dst, __ubuf__ half* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
47 __aicore__ inline void asc_sqrt(__ubuf__ float* dst, __ubuf__ float* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)49 __aicore__ inline void asc_sqrt(__ubuf__ float* dst, __ubuf__ float* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_compare/asc_eq.md+1-0
@@ -33,6 +33,7 @@
33## 函数原型33## 函数原型
34 34 
35- 高维切分计算35- 高维切分计算
36+ 
36 ```cpp37 ```cpp
37 // 比较结果存入寄存器38 // 比较结果存入寄存器
38 __aicore__ inline void asc_eq(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)39 __aicore__ inline void asc_eq(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_compare/asc_ge_scalar.md+1-0
@@ -33,6 +33,7 @@
33## 函数原型33## 函数原型
34 34 
35- 高维切分计算35- 高维切分计算
36+ 
36 ```cpp37 ```cpp
37 __aicore__ inline void asc_ge_scalar(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)38 __aicore__ inline void asc_ge_scalar(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
38 __aicore__ inline void asc_ge_scalar(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)39 __aicore__ inline void asc_ge_scalar(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_compare/asc_gt_scalar.md+1-0
@@ -33,6 +33,7 @@ src中的每个元素逐个与标量value比较大小，如果某个位置上的
33## 函数原型33## 函数原型
34 34 
35- 高维切分计算35- 高维切分计算
36+ 
36 ```cpp37 ```cpp
37 __aicore__ inline void asc_gt_scalar(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)38 __aicore__ inline void asc_gt_scalar(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
38 __aicore__ inline void asc_gt_scalar(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)39 __aicore__ inline void asc_gt_scalar(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_compare/asc_le.md+1-0
@@ -34,6 +34,7 @@
34## 函数原型34## 函数原型
35 35 
36- 高维切分计算36- 高维切分计算
37+ 
37 ```cpp38 ```cpp
38 // 结果输出到寄存器中39 // 结果输出到寄存器中
39 __aicore__ inline void asc_le(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)40 __aicore__ inline void asc_le(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_compare/asc_le_scalar.md+1-0
@@ -33,6 +33,7 @@
33## 函数原型33## 函数原型
34 34 
35- 高维切分计算35- 高维切分计算
36+ 
36 ```cpp37 ```cpp
37 __aicore__ inline void asc_le_scalar(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)38 __aicore__ inline void asc_le_scalar(__ubuf__ uint8_t* dst, __ubuf__ half* src, half value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
38 __aicore__ inline void asc_le_scalar(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)39 __aicore__ inline void asc_le_scalar(__ubuf__ uint8_t* dst, __ubuf__ float* src, float value, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_compare/asc_ne.md+1-0
@@ -34,6 +34,7 @@
34## 函数原型34## 函数原型
35 35 
36- 高维切分计算36- 高维切分计算
37+ 
37 ```cpp38 ```cpp
38 // 结果输出到寄存器中39 // 结果输出到寄存器中
39 __aicore__ inline void asc_ne(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)40 __aicore__ inline void asc_ne(__ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_convert/asc_bfloat162float.md+2-0
@@ -33,11 +33,13 @@
33## 函数原型33## 函数原型
34 34 
35- 前n个数据计算35- 前n个数据计算
36+ 
36```cpp37```cpp
37__aicore__ inline void asc_bfloat162float(__ubuf__ float* dst, __ubuf__ bfloat16_t* src, uint32_t count)38__aicore__ inline void asc_bfloat162float(__ubuf__ float* dst, __ubuf__ bfloat16_t* src, uint32_t count)
38```39```
39 40 
40- 高维切分计算41- 高维切分计算
42+ 
41```cpp43```cpp
42__aicore__ inline void asc_bfloat162float(__ubuf__ float* dst, __ubuf__ bfloat16_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)44__aicore__ inline void asc_bfloat162float(__ubuf__ float* dst, __ubuf__ bfloat16_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
43```45```
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_convert/asc_bfloat162int32.md+2-0
@@ -39,6 +39,7 @@
39## 函数原型39## 函数原型
40 40 
41- 前n个数据计算41- 前n个数据计算
42+ 
42 ```cpp43 ```cpp
43 // RINT舍入模式44 // RINT舍入模式
44 __aicore__ inline void asc_bfloat162int32_rn(__ubuf__ int32_t* dst, __ubuf__ bfloat16_t* src, uint32_t count)45 __aicore__ inline void asc_bfloat162int32_rn(__ubuf__ int32_t* dst, __ubuf__ bfloat16_t* src, uint32_t count)
@@ -57,6 +58,7 @@
57 ```58 ```
58 59 
59- 高维切分计算60- 高维切分计算
61+ 
60 ```cpp62 ```cpp
61 // RINT舍入模式63 // RINT舍入模式
62 __aicore__ inline void asc_bfloat162int32_rn(__ubuf__ int32_t* dst, __ubuf__ bfloat16_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)64 __aicore__ inline void asc_bfloat162int32_rn(__ubuf__ int32_t* dst, __ubuf__ bfloat16_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_convert/asc_float2bfloat16.md+1-0
@@ -39,6 +39,7 @@
39## 函数原型39## 函数原型
40 40 
41- 前n个数据连续计算41- 前n个数据连续计算
42+ 
42 ```cpp43 ```cpp
43 // RINT舍入模式44 // RINT舍入模式
44 __aicore__ inline void asc_float2bfloat16_rn(__ubuf__ bfloat16_t* dst, __ubuf__ float* src, uint32_t count)45 __aicore__ inline void asc_float2bfloat16_rn(__ubuf__ bfloat16_t* dst, __ubuf__ float* src, uint32_t count)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_convert/asc_half2int16.md+2-0
@@ -38,6 +38,7 @@
38## 函数原型38## 函数原型
39 39 
40- 前n个数据计算40- 前n个数据计算
41+ 
41 ```cpp42 ```cpp
42 // RINT舍入模式43 // RINT舍入模式
43 __aicore__ inline void asc_half2int16_rn(__ubuf__ int16_t* dst, __ubuf__ half* src, uint32_t count)44 __aicore__ inline void asc_half2int16_rn(__ubuf__ int16_t* dst, __ubuf__ half* src, uint32_t count)
@@ -52,6 +53,7 @@
52 ```53 ```
53 54 
54- 高维切分计算55- 高维切分计算
56+ 
55 ```cpp57 ```cpp
56 // RINT舍入模式58 // RINT舍入模式
57 __aicore__ inline void asc_half2int16_rn(__ubuf__ int16_t* dst, __ubuf__ half* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)59 __aicore__ inline void asc_half2int16_rn(__ubuf__ int16_t* dst, __ubuf__ half* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_fused/asc_fma.md+2-0
@@ -36,6 +36,7 @@ $$
36## 函数原型36## 函数原型
37 37 
38- 前n个数据计算38- 前n个数据计算
39+ 
39 ```cpp40 ```cpp
40 __aicore__ inline void asc_fma(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)41 __aicore__ inline void asc_fma(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)
41 __aicore__ inline void asc_fma(__ubuf__ float* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)42 __aicore__ inline void asc_fma(__ubuf__ float* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint32_t count)
@@ -43,6 +44,7 @@ $$
43 ```44 ```
44 45 
45- 高维切分计算46- 高维切分计算
47+ 
46 ```cpp48 ```cpp
47 __aicore__ inline void asc_fma(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)49 __aicore__ inline void asc_fma(__ubuf__ half* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
48 __aicore__ inline void asc_fma(__ubuf__ float* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)50 __aicore__ inline void asc_fma(__ubuf__ float* dst, __ubuf__ half* src0, __ubuf__ half* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_fused/asc_relu.md+2-0
@@ -36,6 +36,7 @@ $$
36## 函数原型36## 函数原型
37 37 
38- 前n个数据计算38- 前n个数据计算
39+ 
39 ```cpp40 ```cpp
40 __aicore__ inline void asc_relu(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)41 __aicore__ inline void asc_relu(__ubuf__ half* dst, __ubuf__ half* src, uint32_t count)
41 __aicore__ inline void asc_relu(__ubuf__ int32_t* dst, __ubuf__ int32_t* src, uint32_t count)42 __aicore__ inline void asc_relu(__ubuf__ int32_t* dst, __ubuf__ int32_t* src, uint32_t count)
@@ -43,6 +44,7 @@ $$
43 ```44 ```
44 45 
45- 高维切分计算46- 高维切分计算
47+ 
46 ```cpp48 ```cpp
47 __aicore__ inline void asc_relu(__ubuf__ half* dst, __ubuf__ half* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)49 __aicore__ inline void asc_relu(__ubuf__ half* dst, __ubuf__ half* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
48 __aicore__ inline void asc_relu(__ubuf__ int32_t* dst, __ubuf__ int32_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)50 __aicore__ inline void asc_relu(__ubuf__ int32_t* dst, __ubuf__ int32_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_logic/asc_not.md+2-0
@@ -36,12 +36,14 @@ $$
36## 函数原型36## 函数原型
37 37 
38- 前n个数据计算38- 前n个数据计算
39+ 
39 ```cpp40 ```cpp
40 __aicore__ inline void asc_not(__ubuf__ int16_t* dst, __ubuf__ int16_t* src, uint32_t count)41 __aicore__ inline void asc_not(__ubuf__ int16_t* dst, __ubuf__ int16_t* src, uint32_t count)
41 __aicore__ inline void asc_not(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src, uint32_t count)42 __aicore__ inline void asc_not(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src, uint32_t count)
42 ```43 ```
43 44 
44- 高维切分计算45- 高维切分计算
46+ 
45 ```cpp47 ```cpp
46 __aicore__ inline void asc_not(__ubuf__ int16_t* dst, __ubuf__ int16_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)48 __aicore__ inline void asc_not(__ubuf__ int16_t* dst, __ubuf__ int16_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
47 __aicore__ inline void asc_not(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)49 __aicore__ inline void asc_not(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src, uint8_t repeat, uint16_t dst_block_stride, uint16_t src_block_stride, uint16_t dst_repeat_stride, uint16_t src_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_logic/asc_or.md+2-0
@@ -36,12 +36,14 @@ $$
36## 函数原型36## 函数原型
37 37 
38- 前n个数据计算38- 前n个数据计算
39+ 
39 ```cpp40 ```cpp
40 __aicore__ inline void asc_or(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint32_t count)41 __aicore__ inline void asc_or(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint32_t count)
41 __aicore__ inline void asc_or(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src0, __ubuf__ uint16_t* src1, uint32_t count)42 __aicore__ inline void asc_or(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src0, __ubuf__ uint16_t* src1, uint32_t count)
42 ```43 ```
43 44 
44- 高维切分计算45- 高维切分计算
46+ 
45 ```cpp47 ```cpp
46 __aicore__ inline void asc_or(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)48 __aicore__ inline void asc_or(__ubuf__ int16_t* dst, __ubuf__ int16_t* src0, __ubuf__ int16_t* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
47 __aicore__ inline void asc_or(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src0, __ubuf__ uint16_t* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)49 __aicore__ inline void asc_or(__ubuf__ uint16_t* dst, __ubuf__ uint16_t* src0, __ubuf__ uint16_t* src1, uint8_t repeat, uint8_t dst_block_stride, uint8_t src0_block_stride, uint8_t src1_block_stride, uint8_t dst_repeat_stride, uint8_t src0_repeat_stride, uint8_t src1_repeat_stride)
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_mask_config/asc_set_vector_mask.md+0-1
@@ -184,7 +184,6 @@ PIPE_S
184 return passed ? 0 : 1;184 return passed ? 0 : 1;
185 }185 }
186 ```186 ```
187- 
188<!-- end id18 -->187<!-- end id18 -->
189 188 
190<!-- npu="A3,910b" id19 -->189<!-- npu="A3,910b" id19 -->
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_permute_sel/asc_set_va_reg.md+0-1
@@ -81,7 +81,6 @@ PIPE_V
81```bash81```bash
82bisheng example.asc -o main --npu-arch=dav-3510 && ./main82bisheng example.asc -o main --npu-arch=dav-3510 && ./main
83```83```
84- 
85<!-- end id8 -->84<!-- end id8 -->
86 85 
87```cpp86```cpp
Mdocs/zh/api/SIMD-API/c_api/vector_compute/vector_permute_sel/asc_transto5hd.md+0-1
@@ -106,7 +106,6 @@ PIPE_V
106```bash106```bash
107bisheng example.asc -o main --npu-arch=dav-3510 && ./main107bisheng example.asc -o main --npu-arch=dav-3510 && ./main
108```108```
109- 
110<!-- end id8 -->109<!-- end id8 -->
111 110 
112一次运行会分别验证下面两个场景：111一次运行会分别验证下面两个场景：
Mdocs/zh/api/SIMT-API/memory_access_functions/asc_ldca.md+3-0
@@ -29,10 +29,13 @@
29 29 
30首先从Data Cache加载缓存数据，若未命中，则尝试从L2 Cache加载。L2 Cache与Global Memory之间的数据一致性由硬件保证，但Data Cache和Global Memory之间的数据一致性并不能保证。如果Data Cache和L2 Cache中均未找到所需数据，则从Global Memory中读取数据，然后将其缓存到L2 Cache和Data Cache中。 30首先从Data Cache加载缓存数据，若未命中，则尝试从L2 Cache加载。L2 Cache与Global Memory之间的数据一致性由硬件保证，但Data Cache和Global Memory之间的数据一致性并不能保证。如果Data Cache和L2 Cache中均未找到所需数据，则从Global Memory中读取数据，然后将其缓存到L2 Cache和Data Cache中。
31默认访存的底层实现与本接口一致，例如：31默认访存的底层实现与本接口一致，例如：
32+ 
32```cpp33```cpp
33dst[idx] = src[idx];34dst[idx] = src[idx];
34``` 35```
36+ 
35其底层实现等价于 37其底层实现等价于
38+ 
36```cpp39```cpp
37dst[idx] = asc_ldca(src + idx);40dst[idx] = asc_ldca(src + idx);
38```41```
Mdocs/zh/api/Utils-API/tuning_interface/printf.md+2-0
@@ -94,11 +94,13 @@ static __attribute__((noinline)) void printf(const char* fmt, Args&&... args);
94- SIMT中使用printf接口会增加算子运行时间，请控制打印数据量，避免算子超时。建议在代码中判断线程ID，仅在部分线程中打印调试信息，减少重复内容的打印，更有利于调试。94- SIMT中使用printf接口会增加算子运行时间，请控制打印数据量，避免算子超时。建议在代码中判断线程ID，仅在部分线程中打印调试信息，减少重复内容的打印，更有利于调试。
95- SIMD场景下，单次调用本接口打印的数据总量不可超过打印大小限制，默认为30KB。使用时应注意，如果超出这个限制，则数据不会被打印。您可以通过[aclInit接口](https://gitcode.com/cann/runtime/blob/master/docs/zh/api_ref/02_initialization_and_deinitialization.md#aclinit)中的"simd\_printf\_fifo\_size\_per\_core"字段进行配置，配置范围为`[1KB,64MB]`。当打印数据量较大时，建议增加缓存空间。pytorch调用和算子入图场景暂不支持该配置。95- SIMD场景下，单次调用本接口打印的数据总量不可超过打印大小限制，默认为30KB。使用时应注意，如果超出这个限制，则数据不会被打印。您可以通过[aclInit接口](https://gitcode.com/cann/runtime/blob/master/docs/zh/api_ref/02_initialization_and_deinitialization.md#aclinit)中的"simd\_printf\_fifo\_size\_per\_core"字段进行配置，配置范围为`[1KB,64MB]`。当打印数据量较大时，建议增加缓存空间。pytorch调用和算子入图场景暂不支持该配置。
96- 如果需要使用"AscendC::printf"形式的打印，则必须额外包含"kernel\_operator.h"头文件，示例如下。96- 如果需要使用"AscendC::printf"形式的打印，则必须额外包含"kernel\_operator.h"头文件，示例如下。
97+ 
97 ```cpp98 ```cpp
98 #include "kernel_operator.h"99 #include "kernel_operator.h"
99 100 
100 AscendC::printf("fmt string %d\n", 0x123);101 AscendC::printf("fmt string %d\n", 0x123);
101 ```102 ```
103+ 
102- 由于在simd vf中不能直接访问gm地址，因此使用simd vf中的printf时，需要先定义一个ubuf的字符串，该字符串即为原来的fmt变量。104- 由于在simd vf中不能直接访问gm地址，因此使用simd vf中的printf时，需要先定义一个ubuf的字符串，该字符串即为原来的fmt变量。
103 105 
104 ```cpp106 ```cpp
Mdocs/zh/guide/figures/image_compression_guide.md+4-0
@@ -18,6 +18,7 @@
18| 6 | addr_alloc_aft.png | **125.12 KB** | < 80 KB | 需压缩36% |18| 6 | addr_alloc_aft.png | **125.12 KB** | < 80 KB | 需压缩36% |
19 19 
20**图片文件位置**：20**图片文件位置**：
21+ 
21```22```
22C:\work_for_code\doc_design\asc-devkit\docs\guide\figures\23C:\work_for_code\doc_design\asc-devkit\docs\guide\figures\
23```24```
@@ -133,6 +134,7 @@ optipng -o7 addr_alloc_aft.png
133### 操作步骤截图说明134### 操作步骤截图说明
134 135 
135**步骤1**：打开 TinyPNG 网站136**步骤1**：打开 TinyPNG 网站
137+ 
136```138```
137https://tinypng.com/139https://tinypng.com/
138```140```
@@ -151,9 +153,11 @@ https://tinypng.com/
151 153 
152**步骤5**：替换原文件154**步骤5**：替换原文件
153- 将压缩后的图片复制到：155- 将压缩后的图片复制到：
156+ 
154 ```157 ```
155 C:\work_for_code\doc_design\asc-devkit\docs\guide\figures\158 C:\work_for_code\doc_design\asc-devkit\docs\guide\figures\
156 ```159 ```
160+ 
157- 替换原有的大图片文件161- 替换原有的大图片文件
158 162 
159---163---
Mdocs/zh/guide/getting_started/quick_start/simd_programming/add_op_quick_start.md+8-0
@@ -37,6 +37,7 @@
37 - **Device端核函数（Kernel）实现**：37 - **Device端核函数（Kernel）实现**：
38 Device端部分示例如下：38 Device端部分示例如下：
39 - **基于C语言API实现Memory矢量计算示例**39 - **基于C语言API实现Memory矢量计算示例**
40+ 
40 ```cpp41 ```cpp
41 __vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z)42 __vector__ __global__ void add_custom(__gm__ float* x, __gm__ float* y, __gm__ float* z)
42 {43 {
@@ -68,6 +69,7 @@
68 asc_sync();69 asc_sync();
69 }70 }
70 ```71 ```
72+ 
71 > [!NOTE]说明73 > [!NOTE]说明
72 > - 本Memory矢量计算示例支持以下型号：74 > - 本Memory矢量计算示例支持以下型号：
73 > <!-- npu="A3" id1 -->75 > <!-- npu="A3" id1 -->
@@ -113,6 +115,7 @@
113 AscendC::PipeBarrier<PIPE_ALL>();115 AscendC::PipeBarrier<PIPE_ALL>();
114 }116 }
115 ```117 ```
118+ 
116 > [!NOTE]说明119 > [!NOTE]说明
117 > - 该样例支持以下型号：120 > - 该样例支持以下型号：
118 > <!-- npu="950" id3 -->121 > <!-- npu="950" id3 -->
@@ -129,6 +132,7 @@
129 - **Host端代码实现**：132 - **Host端代码实现**：
130 133 
131 Host端通过`<<<>>>`语法糖调用Device端核函数（Kernel），示例代码片段如下：134 Host端通过`<<<>>>`语法糖调用Device端核函数（Kernel），示例代码片段如下：
135+ 
132 ```cpp136 ```cpp
133 int32_t main(int argc, char const *argv[])137 int32_t main(int argc, char const *argv[])
134 {138 {
@@ -169,6 +173,7 @@
169- **算子编译与运行**：173- **算子编译与运行**：
170 174 
171 CMake配置文件示例：175 CMake配置文件示例：
176+ 
172 ```cmake177 ```cmake
173 cmake_minimum_required(VERSION 3.16)178 cmake_minimum_required(VERSION 3.16)
174 179 
@@ -190,12 +195,15 @@
190 $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>195 $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>
191 )196 )
192 ```197 ```
198+ 
193 编译与执行示例：199 编译与执行示例：
200+ 
194 ```bash201 ```bash
195 mkdir -p build && cd build; # 创建并进入build目录202 mkdir -p build && cd build; # 创建并进入build目录
196 cmake ..;make -j; # 编译工程203 cmake ..;make -j; # 编译工程
197 ./c_api_add_example # 运行样例204 ./c_api_add_example # 运行样例
198 ```205 ```
206+ 
199 > [!NOTE]说明207 > [!NOTE]说明
200 > - 编译选项`--npu-arch`用于指定NPU架构版本，`dav-`后面的数字为架构版本号，请替换为您实际使用的版本。各AI处理器型号与架构版本的对应关系请查阅[AI处理器型号和 \_\_NPU_ARCH\_\_ 的对应关系](../../../programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。208 > - 编译选项`--npu-arch`用于指定NPU架构版本，`dav-`后面的数字为架构版本号，请替换为您实际使用的版本。各AI处理器型号与架构版本的对应关系请查阅[AI处理器型号和 \_\_NPU_ARCH\_\_ 的对应关系](../../../programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
201 209 
Mdocs/zh/guide/getting_started/quick_start/simd_programming/hello_world.md+4-0
@@ -11,18 +11,21 @@
11- **Device端核函数（Kernel）实现**：11- **Device端核函数（Kernel）实现**：
12 12 
13 后缀名为`*.asc`的代码文件包含Host端与Device端代码，其Device端部分示例如下：13 后缀名为`*.asc`的代码文件包含Host端与Device端代码，其Device端部分示例如下：
14+ 
14 ```cpp15 ```cpp
15 __global__ __vector__ void hello_world()16 __global__ __vector__ void hello_world()
16 {17 {
17 printf("Hello World!!!\n");18 printf("Hello World!!!\n");
18 }19 }
19 ```20 ```
21+ 
20 > [!NOTE]说明22 > [!NOTE]说明
21 > - SIMD算子的核函数（Kernel）需要额外的修饰符，如[`__vector__`](../../../programming_guide/language_extension/simd_builtin_keywords.md)修饰符说明该算子仅在向量计算单元上执行。23 > - SIMD算子的核函数（Kernel）需要额外的修饰符，如[`__vector__`](../../../programming_guide/language_extension/simd_builtin_keywords.md)修饰符说明该算子仅在向量计算单元上执行。
22 24
23- **Host端代码实现**：25- **Host端代码实现**：
24 26
25 Host端通过<<<>>>语法糖调用Device端代码片段。27 Host端通过<<<>>>语法糖调用Device端代码片段。
28+ 
26 ```cpp29 ```cpp
27 int main(int argc, char const* argv[])30 int main(int argc, char const* argv[])
28 {31 {
@@ -47,6 +50,7 @@
47 bisheng hello_world.asc --npu-arch=dav-2201 -o demo50 bisheng hello_world.asc --npu-arch=dav-2201 -o demo
48 ./demo51 ./demo
49 ```52 ```
53+ 
50 执行本样例，将打印了核号和`Hello World!!!`信息。54 执行本样例，将打印了核号和`Hello World!!!`信息。
51 55 
52 > [!NOTE]说明56 > [!NOTE]说明