已合并
C API系统变量接口资料修改 #5620
C API系统变量接口资料修改 #5620
已合并
qwerwjy创建于 14 天前
125 个文件变更+1248-356
@@ -799,6 +799,7 @@
799 - [asc_get_sub_block_num](SIMD-API/c_api/utils/sys_var/asc_get_sub_block_num.md)799 - [asc_get_sub_block_num](SIMD-API/c_api/utils/sys_var/asc_get_sub_block_num.md)
800 - [asc_get_sys_virtual_base](SIMD-API/c_api/utils/sys_var/asc_get_sys_virtual_base.md)800 - [asc_get_sys_virtual_base](SIMD-API/c_api/utils/sys_var/asc_get_sys_virtual_base.md)
801 - [asc_get_system_cycle](SIMD-API/c_api/utils/sys_var/asc_get_system_cycle.md)801 - [asc_get_system_cycle](SIMD-API/c_api/utils/sys_var/asc_get_system_cycle.md)
802+ - [asc_get_system_clock](SIMD-API/c_api/utils/sys_var/asc_get_system_clock.md)
802 - [asc_get_vf_len](SIMD-API/c_api/utils/sys_var/asc_get_vf_len.md)803 - [asc_get_vf_len](SIMD-API/c_api/utils/sys_var/asc_get_vf_len.md)
803 - [asc_set_ffts_base_addr](SIMD-API/c_api/utils/sys_var/asc_set_ffts_base_addr.md)804 - [asc_set_ffts_base_addr](SIMD-API/c_api/utils/sys_var/asc_set_ffts_base_addr.md)
804 - [矩阵计算](SIMD-API/c_api/cube_compute/cube_compute.md)805 - [矩阵计算](SIMD-API/c_api/cube_compute/cube_compute.md)
@@ -89,7 +89,6 @@ PIPE_S
89将以下代码保存为`example.asc`后,执行对应的编译运行命令。89将以下代码保存为`example.asc`后,执行对应的编译运行命令。
90 90 
91<!-- npu="950" id8 -->91<!-- npu="950" id8 -->
92- 
93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:92以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
94 93 
95```bash94```bash
@@ -76,7 +76,6 @@ PIPE_S
76将以下代码保存为`example.asc`后,执行对应的编译运行命令。76将以下代码保存为`example.asc`后,执行对应的编译运行命令。
77 77 
78<!-- npu="950" id8 -->78<!-- npu="950" id8 -->
79- 
80以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:79以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
81 80 
82```bash81```bash
@@ -85,7 +85,6 @@ PIPE_S
85将以下代码保存为`example.asc`后,执行对应的编译运行命令。85将以下代码保存为`example.asc`后,执行对应的编译运行命令。
86 86 
87<!-- npu="950" id8 -->87<!-- npu="950" id8 -->
88- 
89以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:88以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
90 89 
91```bash90```bash
@@ -82,7 +82,6 @@ PIPE_S
82将以下代码保存为`example.asc`后,执行对应的编译运行命令。82将以下代码保存为`example.asc`后,执行对应的编译运行命令。
83 83 
84<!-- npu="950" id8 -->84<!-- npu="950" id8 -->
85- 
86以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:85以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
87 86 
88```bash87```bash
@@ -76,7 +76,6 @@ PIPE_S
76将以下代码保存为`example.asc`后,执行对应的编译运行命令。76将以下代码保存为`example.asc`后,执行对应的编译运行命令。
77 77 
78<!-- npu="950" id8 -->78<!-- npu="950" id8 -->
79- 
80以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:79以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
81 80 
82```bash81```bash
@@ -82,7 +82,6 @@ PIPE_S
82将以下代码保存为`example.asc`后,执行对应的编译运行命令。82将以下代码保存为`example.asc`后,执行对应的编译运行命令。
83 83 
84<!-- npu="950" id8 -->84<!-- npu="950" id8 -->
85- 
86以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:85以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
87 86 
88```bash87```bash
@@ -66,6 +66,7 @@ C API文档按最细粒度公共头文件归类。除试验接口外,整体使
66|[asc_get_sub_block_num](utils/sys_var/asc_get_sub_block_num.md)|分离模式下,获取一个AI Core上Cube Core(AIC)或者Vector Core(AIV)的数量。|`c_api/utils/sys_var.h`|66|[asc_get_sub_block_num](utils/sys_var/asc_get_sub_block_num.md)|分离模式下,获取一个AI Core上Cube Core(AIC)或者Vector Core(AIV)的数量。|`c_api/utils/sys_var.h`|
67|[asc_get_sys_virtual_base](utils/sys_var/asc_get_sys_virtual_base.md)|获取系统虚拟基地址。|`c_api/utils/sys_var.h`|67|[asc_get_sys_virtual_base](utils/sys_var/asc_get_sys_virtual_base.md)|获取系统虚拟基地址。|`c_api/utils/sys_var.h`|
68|[asc_get_system_cycle](utils/sys_var/asc_get_system_cycle.md)|获取当前系统cycle数。|`c_api/utils/sys_var.h`|68|[asc_get_system_cycle](utils/sys_var/asc_get_system_cycle.md)|获取当前系统cycle数。|`c_api/utils/sys_var.h`|
69+|[asc_get_system_clock](utils/sys_var/asc_get_system_clock.md)|获取当前系统时钟时间值(微秒)。|`c_api/utils/sys_var.h`|
69|[asc_get_vf_len](utils/sys_var/asc_get_vf_len.md)|获取Tensor位宽VL(Vector Length)的大小。|`c_api/utils/sys_var.h`|70|[asc_get_vf_len](utils/sys_var/asc_get_vf_len.md)|获取Tensor位宽VL(Vector Length)的大小。|`c_api/utils/sys_var.h`|
70|[asc_set_ffts_base_addr](utils/sys_var/asc_set_ffts_base_addr.md)|在[asc_sync_block_arrive](sync/asc_sync_block_arrive.md)和[asc_sync_block_wait](sync/asc_sync_block_wait.md)之前使用,设置核间同步寄存器的基地址。|`c_api/utils/sys_var.h`|71|[asc_set_ffts_base_addr](utils/sys_var/asc_set_ffts_base_addr.md)|在[asc_sync_block_arrive](sync/asc_sync_block_arrive.md)和[asc_sync_block_wait](sync/asc_sync_block_wait.md)之前使用,设置核间同步寄存器的基地址。|`c_api/utils/sys_var.h`|
71 72 
@@ -56,16 +56,17 @@ PIPE_S
56- 本接口需在矩阵乘加指令([asc_mmad](./asc_mmad.md))执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。56- 本接口需在矩阵乘加指令([asc_mmad](./asc_mmad.md))执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。
57- 与[asc_enable_hf32](./asc_enable_hf32.md)作用相反,二者不同时生效。57- 与[asc_enable_hf32](./asc_enable_hf32.md)作用相反,二者不同时生效。
58 58 
59-<!-- npu="950" id8 -->
60## 调用示例59## 调用示例
61 60 
62将代码保存为`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)。
63 62 
63+<!-- npu="950" id8 -->
64以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:64以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
65 65 
66```bash66```bash
67bisheng example.asc -o main --npu-arch=dav-3510 && ./main67bisheng example.asc -o main --npu-arch=dav-3510 && ./main
68```68```
69+<!-- end id8 -->
69 70 
70以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。71以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
71 72 
@@ -159,4 +160,3 @@ int main()
159 return passed ? 0 : 1;160 return passed ? 0 : 1;
160}161}
161```162```
162-<!-- end id8 -->
@@ -56,16 +56,17 @@ PIPE_S
56- 需在[asc_mmad](./asc_mmad.md)执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。56- 需在[asc_mmad](./asc_mmad.md)执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。
57- 与[asc_enable_hif8](./asc_enable_hif8.md)作用相反,二者不同时生效。57- 与[asc_enable_hif8](./asc_enable_hif8.md)作用相反,二者不同时生效。
58 58 
59-<!-- npu="950" id8 -->
60## 调用示例59## 调用示例
61 60 
62将代码保存为`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)。
63 62 
63+<!-- npu="950" id8 -->
64以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:64以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
65 65 
66```bash66```bash
67bisheng example.asc -o main --npu-arch=dav-3510 && ./main67bisheng example.asc -o main --npu-arch=dav-3510 && ./main
68```68```
69+<!-- end id8 -->
69 70 
70以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。71以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
71 72 
@@ -175,4 +176,3 @@ int main()
175 return passed ? 0 : 1;176 return passed ? 0 : 1;
176}177}
177```178```
178-<!-- end id8 -->
@@ -112,16 +112,17 @@ PIPE_S
112- HF32模式启用后,FP32转换为HF32格式的舍入模式由配套的[asc_set_hf32_round_mode](asc_set_hf32_round_mode.md)接口配置,须在本接口之后调用该配套接口配置舍入模式,否则舍入模式沿用上次配置或默认值。112- HF32模式启用后,FP32转换为HF32格式的舍入模式由配套的[asc_set_hf32_round_mode](asc_set_hf32_round_mode.md)接口配置,须在本接口之后调用该配套接口配置舍入模式,否则舍入模式沿用上次配置或默认值。
113- HF32模式仅对输入矩阵数据类型为`float`的矩阵乘加运算场景生效。113- HF32模式仅对输入矩阵数据类型为`float`的矩阵乘加运算场景生效。
114 114 
115-<!-- npu="950" id12 -->
116## 调用示例115## 调用示例
117 116 
118将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。117将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
119 118 
119+<!-- npu="950" id12 -->
120以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:120以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
121 121 
122```bash122```bash
123bisheng example.asc -o main --npu-arch=dav-3510 && ./main123bisheng example.asc -o main --npu-arch=dav-3510 && ./main
124```124```
125+<!-- end id12 -->
125 126 
126以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。127以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
127 128 
@@ -230,4 +231,3 @@ int main()
230 return passed ? 0 : 1;231 return passed ? 0 : 1;
231}232}
232```233```
233-<!-- end id12 -->
@@ -58,16 +58,17 @@ PIPE_S
58- 本接口配置的舍入模式仅在HF32模式开启期间生效,需先开启HF32模式再调用本接口,否则舍入模式配置不产生实际作用。58- 本接口配置的舍入模式仅在HF32模式开启期间生效,需先开启HF32模式再调用本接口,否则舍入模式配置不产生实际作用。
59- 舍入模式配置后会持续生效,HF32模式关闭后再次开启仍将沿用上次的舍入模式配置,如需切换舍入模式,请重新调用本接口。59- 舍入模式配置后会持续生效,HF32模式关闭后再次开启仍将沿用上次的舍入模式配置,如需切换舍入模式,请重新调用本接口。
60 60 
61-<!-- npu="950" id8 -->
62## 调用示例61## 调用示例
63 62 
64将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。63将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
65 64 
65+<!-- npu="950" id8 -->
66以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:66以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
67 67 
68```bash68```bash
69bisheng example.asc -o main --npu-arch=dav-3510 && ./main69bisheng example.asc -o main --npu-arch=dav-3510 && ./main
70```70```
71+<!-- end id8 -->
71 72 
72以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。73以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
73 74 
@@ -157,4 +158,3 @@ int main()
157 return passed ? 0 : 1;158 return passed ? 0 : 1;
158}159}
159```160```
160-<!-- end id8 -->
@@ -262,16 +262,17 @@ PIPE_M
262 262 
263 注意,应避免nan输入,否则可能会产生执行报错;整数类型仅支持饱和模式。263 注意,应避免nan输入,否则可能会产生执行报错;整数类型仅支持饱和模式。
264 264 
265-<!-- npu="950" id18 -->
266## 调用示例265## 调用示例
267 266 
268将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。267将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
269 268 
269+<!-- npu="950" id18 -->
270以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:270以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
271 271 
272```bash272```bash
273bisheng example.asc -o main --npu-arch=dav-3510 && ./main273bisheng example.asc -o main --npu-arch=dav-3510 && ./main
274```274```
275+<!-- end id18 -->
275 276 
276以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。277以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
277 278 
@@ -380,4 +381,3 @@ int main()
380```381```
381 382 
382更多场景使用样例请参考[asc_mmad样例](../../../../../../examples/02_simd_c_api/03_c_api/03_matrix_compute/mmad)。383更多场景使用样例请参考[asc_mmad样例](../../../../../../examples/02_simd_c_api/03_c_api/03_matrix_compute/mmad)。
383-<!-- end id18 -->
@@ -189,16 +189,17 @@ PIPE_M
189 189 
190 边界值计算结果请见[MmadMx接口边界值汇总](../../../appendix/mmad_mx_interface_boundary_value_summary.md)。190 边界值计算结果请见[MmadMx接口边界值汇总](../../../appendix/mmad_mx_interface_boundary_value_summary.md)。
191 191 
192-<!-- npu="950" id8 -->
193## 调用示例192## 调用示例
194 193 
195将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。194将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
196 195 
196+<!-- npu="950" id8 -->
197以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:197以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
198 198 
199```bash199```bash
200bisheng example.asc -o main --npu-arch=dav-3510 && ./main200bisheng example.asc -o main --npu-arch=dav-3510 && ./main
201```201```
202+<!-- end id8 -->
202 203 
203```cpp204```cpp
204#include <cmath>205#include <cmath>
@@ -355,4 +356,3 @@ int main()
355```356```
356 357 
357更多场景使用样例请参考[asc_mmad_mx样例](../../../../../../examples/02_simd_c_api/03_c_api/03_matrix_compute/mmad_mx)。358更多场景使用样例请参考[asc_mmad_mx样例](../../../../../../examples/02_simd_c_api/03_c_api/03_matrix_compute/mmad_mx)。
358-<!-- end id8 -->
@@ -53,16 +53,17 @@ PIPE_S
53- 本接口需在矩阵乘加指令([asc_mmad](asc_mmad.md))执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。53- 本接口需在矩阵乘加指令([asc_mmad](asc_mmad.md))执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。
54- FP32模式启用后会持续生效,不会自动关闭。后续矩阵乘加指令若不显式重新配置,将沿用当前模式。如需开启HF32模式,请重新调用[asc_enable_hf32](asc_enable_hf32.md)接口。54- FP32模式启用后会持续生效,不会自动关闭。后续矩阵乘加指令若不显式重新配置,将沿用当前模式。如需开启HF32模式,请重新调用[asc_enable_hf32](asc_enable_hf32.md)接口。
55 55 
56-<!-- npu="950" id8 -->
57## 调用示例56## 调用示例
58 57 
59将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。58将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
60 59 
60+<!-- npu="950" id8 -->
61以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:61以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
62 62 
63```bash63```bash
64bisheng example.asc -o main --npu-arch=dav-3510 && ./main64bisheng example.asc -o main --npu-arch=dav-3510 && ./main
65```65```
66+<!-- end id8 -->
66 67 
67以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。68以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
68 69 
@@ -156,4 +157,3 @@ int main()
156 return passed ? 0 : 1;157 return passed ? 0 : 1;
157}158}
158```159```
159-<!-- end id8 -->
@@ -57,16 +57,17 @@ PIPE_S
57- 本接口配置的舍入模式仅在HF32模式开启期间生效,需先开启HF32模式再调用本接口,否则舍入模式配置不产生实际作用。57- 本接口配置的舍入模式仅在HF32模式开启期间生效,需先开启HF32模式再调用本接口,否则舍入模式配置不产生实际作用。
58- 舍入模式配置后会持续生效,HF32模式关闭后再次开启仍将沿用上次的舍入模式配置,如需切换舍入模式,请重新调用本接口。58- 舍入模式配置后会持续生效,HF32模式关闭后再次开启仍将沿用上次的舍入模式配置,如需切换舍入模式,请重新调用本接口。
59 59 
60-<!-- npu="950" id8 -->
61## 调用示例60## 调用示例
62 61 
63将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。62将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
64 63 
64+<!-- npu="950" id8 -->
65以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:65以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
66 66 
67```bash67```bash
68bisheng example.asc -o main --npu-arch=dav-3510 && ./main68bisheng example.asc -o main --npu-arch=dav-3510 && ./main
69```69```
70+<!-- end id8 -->
70 71 
71以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。72以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
72 73 
@@ -157,4 +158,3 @@ int main()
157 return passed ? 0 : 1;158 return passed ? 0 : 1;
158}159}
159```160```
160-<!-- end id8 -->
@@ -59,16 +59,17 @@ PIPE_S
59- 本接口需在矩阵乘加指令([asc_mmad](asc_mmad.md)、[asc_mmad_mx](asc_mmad_mx.md))执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。59- 本接口需在矩阵乘加指令([asc_mmad](asc_mmad.md)、[asc_mmad_mx](asc_mmad_mx.md))执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。
60- 方向配置一旦写入会持续生效,后续矩阵乘加指令若不显式重新配置,将沿用当前方向配置。如需切换为逐行生成矩阵计算结果分形,请重新调用[asc_set_mmad_direction_n](asc_set_mmad_direction_n.md)接口。60- 方向配置一旦写入会持续生效,后续矩阵乘加指令若不显式重新配置,将沿用当前方向配置。如需切换为逐行生成矩阵计算结果分形,请重新调用[asc_set_mmad_direction_n](asc_set_mmad_direction_n.md)接口。
61 61 
62-<!-- npu="950" id8 -->
63## 调用示例62## 调用示例
64 63 
65将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。64将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
66 65 
66+<!-- npu="950" id8 -->
67以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:67以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
68 68 
69```bash69```bash
70bisheng example.asc -o main --npu-arch=dav-3510 && ./main70bisheng example.asc -o main --npu-arch=dav-3510 && ./main
71```71```
72+<!-- end id8 -->
72 73 
73以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。74以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
74 75 
@@ -168,4 +169,3 @@ int main()
168 return passed ? 0 : 1;169 return passed ? 0 : 1;
169}170}
170```171```
171-<!-- end id8 -->
@@ -58,16 +58,17 @@ PIPE_S
58- 本接口需在矩阵乘加指令([asc_mmad](asc_mmad.md)、[asc_mmad_mx](asc_mmad_mx.md))执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。58- 本接口需在矩阵乘加指令([asc_mmad](asc_mmad.md)、[asc_mmad_mx](asc_mmad_mx.md))执行前调用,以此来确保模式配置在矩阵乘加计算过程中生效。
59- 方向配置一旦写入会持续生效,后续矩阵乘加指令若不显式重新配置,将沿用当前方向配置。如需切换为逐列生成矩阵计算结果分形,请重新调用[asc_set_mmad_direction_m](asc_set_mmad_direction_m.md)接口。59- 方向配置一旦写入会持续生效,后续矩阵乘加指令若不显式重新配置,将沿用当前方向配置。如需切换为逐列生成矩阵计算结果分形,请重新调用[asc_set_mmad_direction_m](asc_set_mmad_direction_m.md)接口。
60 60 
61-<!-- npu="950" id8 -->
62## 调用示例61## 调用示例
63 62 
64将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。63将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
65 64 
65+<!-- npu="950" id8 -->
66以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:66以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
67 67 
68```bash68```bash
69bisheng example.asc -o main --npu-arch=dav-3510 && ./main69bisheng example.asc -o main --npu-arch=dav-3510 && ./main
70```70```
71+<!-- end id8 -->
71 72 
72以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。73以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
73 74 
@@ -157,4 +158,3 @@ int main()
157 return passed ? 0 : 1;158 return passed ? 0 : 1;
158}159}
159```160```
160-<!-- end id8 -->
@@ -89,16 +89,17 @@ PIPE_S
89- 搬运至L1 Buffer的数据不能重叠。如果步长配置导致重叠写入,硬件不会产生告警或错误,并且不保证重叠数据的写入顺序。89- 搬运至L1 Buffer的数据不能重叠。如果步长配置导致重叠写入,硬件不会产生告警或错误,并且不保证重叠数据的写入顺序。
90- `config`的四个字段均为16bit,接口会将`config`整体写入,不会单独检查字段范围。拼装前需确保各字段满足[表2](#table2)中的取值要求,不得将超出16bit的值直接左移,否则高位会写入相邻字段。90- `config`的四个字段均为16bit,接口会将`config`整体写入,不会单独检查字段范围。拼装前需确保各字段满足[表2](#table2)中的取值要求,不得将超出16bit的值直接左移,否则高位会写入相邻字段。
91 91 
92-<!-- npu="950" id8 -->
93## 调用示例92## 调用示例
94 93 
95将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。94将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
96 95 
96+<!-- npu="950" id8 -->
97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
98 98 
99```bash99```bash
100bisheng example.asc -o main --npu-arch=dav-3510 && ./main100bisheng example.asc -o main --npu-arch=dav-3510 && ./main
101```101```
102+<!-- end id8 -->
102 103 
103以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。104以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
104 105 
@@ -203,4 +204,3 @@ int main()
203 return passed ? 0 : 1;204 return passed ? 0 : 1;
204}205}
205```206```
206-<!-- end id8 -->
@@ -68,16 +68,17 @@ PIPE_S
68- `config[63:48]`不能为0,除`config[63:48]`外的其他bit位不生效,设置为0即可;接口会将`config`整体写入,不会检查字段值。68- `config[63:48]`不能为0,除`config[63:48]`外的其他bit位不生效,设置为0即可;接口会将`config`整体写入,不会检查字段值。
69- 开启Nz2DN功能时,如果`config[63:48]`不等于1,则不能同时开启unitFlag。69- 开启Nz2DN功能时,如果`config[63:48]`不等于1,则不能同时开启unitFlag。
70 70 
71-<!-- npu="950" id8 -->
72## 调用示例71## 调用示例
73 72 
74将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。73将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
75 74 
75+<!-- npu="950" id8 -->
76以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:76以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
77 77 
78```bash78```bash
79bisheng example.asc -o main --npu-arch=dav-3510 && ./main79bisheng example.asc -o main --npu-arch=dav-3510 && ./main
80```80```
81+<!-- end id8 -->
81 82 
82以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。83以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
83 84 
@@ -182,4 +183,3 @@ int main()
182 return passed ? 0 : 1;183 return passed ? 0 : 1;
183}184}
184```185```
185-<!-- end id8 -->
@@ -63,16 +63,17 @@ PIPE_S
63- 本接口仅在AIC上生效,在AIV上调用将直接返回。63- 本接口仅在AIC上生效,在AIV上调用将直接返回。
64- 调用L0C Buffer搬出接口[asc_copy_l0c2gm](asc_copy_l0c2gm/asc_copy_l0c2gm_arch_3510.md)、[asc_copy_l0c2l1](asc_copy_l0c2l1/asc_copy_l0c2l1_arch_3510.md)、[asc_copy_l0c2ub](asc_copy_l0c2ub.md)并且需要进行随路Nz转ND格式之前,必须先调用本接口配置格式转换参数。64- 调用L0C Buffer搬出接口[asc_copy_l0c2gm](asc_copy_l0c2gm/asc_copy_l0c2gm_arch_3510.md)、[asc_copy_l0c2l1](asc_copy_l0c2l1/asc_copy_l0c2l1_arch_3510.md)、[asc_copy_l0c2ub](asc_copy_l0c2ub.md)并且需要进行随路Nz转ND格式之前,必须先调用本接口配置格式转换参数。
65 65 
66-<!-- npu="950" id8 -->
67## 调用示例66## 调用示例
68 67 
69将代码保存为`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)。
70 69 
70+<!-- npu="950" id8 -->
71以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:71以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
72 72 
73```bash73```bash
74bisheng example.asc -o main --npu-arch=dav-3510 && ./main74bisheng example.asc -o main --npu-arch=dav-3510 && ./main
75```75```
76+<!-- end id8 -->
76 77 
77以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。78以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
78 79 
@@ -156,4 +157,3 @@ int main()
156 return passed ? 0 : 1;157 return passed ? 0 : 1;
157}158}
158```159```
159-<!-- end id8 -->
@@ -61,16 +61,17 @@ PIPE_S
61- 开启Nz2DN功能,`src_nz_fractal_stride`不能为0。61- 开启Nz2DN功能,`src_nz_fractal_stride`不能为0。
62- 开启Nz2DN功能时,如果`src_nz_fractal_stride`不等于1,则不能同时开启UnitFlag。62- 开启Nz2DN功能时,如果`src_nz_fractal_stride`不等于1,则不能同时开启UnitFlag。
63 63 
64-<!-- npu="950" id8 -->
65## 调用示例64## 调用示例
66 65 
67将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。66将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
68 67 
68+<!-- npu="950" id8 -->
69以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:69以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
70 70 
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+<!-- end id8 -->
74 75 
75以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。76以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
76 77 
@@ -175,4 +176,3 @@ int main()
175 return passed ? 0 : 1;176 return passed ? 0 : 1;
176}177}
177```178```
178-<!-- end id8 -->
@@ -63,16 +63,17 @@ PIPE_S
63- 本接口非AIC调用直接返回。63- 本接口非AIC调用直接返回。
64- 调用L0C Buffer搬出接口[asc_copy_l0c2gm](../cube_datamove/asc_copy_l0c2gm/asc_copy_l0c2gm_arch_3510.md)、[asc_copy_l0c2l1](../cube_datamove/asc_copy_l0c2l1/asc_copy_l0c2l1_arch_3510.md)、[asc_copy_l0c2ub](../cube_datamove/asc_copy_l0c2ub.md)并且需要进行随路Nz转ND/DN格式之前,即搬出接口内参数`enable_nz2nd``enable_nz2dn`其一设置为`true`时,必须先调用本接口配置格式转换参数。64- 调用L0C Buffer搬出接口[asc_copy_l0c2gm](../cube_datamove/asc_copy_l0c2gm/asc_copy_l0c2gm_arch_3510.md)、[asc_copy_l0c2l1](../cube_datamove/asc_copy_l0c2l1/asc_copy_l0c2l1_arch_3510.md)、[asc_copy_l0c2ub](../cube_datamove/asc_copy_l0c2ub.md)并且需要进行随路Nz转ND/DN格式之前,即搬出接口内参数`enable_nz2nd``enable_nz2dn`其一设置为`true`时,必须先调用本接口配置格式转换参数。
65 65 
66-<!-- npu="950" id8 -->
67## 调用示例66## 调用示例
68 67 
69将代码保存为`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)。
70 69 
70+<!-- npu="950" id8 -->
71以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:71以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
72 72 
73```bash73```bash
74bisheng example.asc -o main --npu-arch=dav-3510 && ./main74bisheng example.asc -o main --npu-arch=dav-3510 && ./main
75```75```
76+<!-- end id8 -->
76 77 
77以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。78以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
78 79 
@@ -156,4 +157,3 @@ int main()
156 return passed ? 0 : 1;157 return passed ? 0 : 1;
157}158}
158```159```
159-<!-- end id8 -->
@@ -92,16 +92,17 @@ PIPE_S
92- 本接口非AIC调用直接返回。92- 本接口非AIC调用直接返回。
93- 本接口需在对应的L0C Buffer搬出接口执行前调用,并在搬出接口中配置对应的Scalar量化模式。93- 本接口需在对应的L0C Buffer搬出接口执行前调用,并在搬出接口中配置对应的Scalar量化模式。
94 94 
95-<!-- npu="950" id8 -->
96## 调用示例95## 调用示例
97 96 
98将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。97将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
99 98 
99+<!-- npu="950" id8 -->
100以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:100以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
101 101 
102```bash102```bash
103bisheng example.asc -o main --npu-arch=dav-3510 && ./main103bisheng example.asc -o main --npu-arch=dav-3510 && ./main
104```104```
105+<!-- end id8 -->
105 106 
106以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。107以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
107 108 
@@ -192,4 +193,3 @@ int main()
192 return passed ? 0 : 1;193 return passed ? 0 : 1;
193}194}
194```195```
195-<!-- end id8 -->
@@ -90,7 +90,6 @@ __simd_callee__ inline void asc_load(vector_float& dst,
90将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。90将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
91 91 
92<!-- npu="950" id8 -->92<!-- npu="950" id8 -->
93- 
94以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
95 94 
96```bash95```bash
@@ -161,7 +161,6 @@ __simd_callee__ inline void asc_loadalign_brc_datablock(vector_half& dst,
161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
162 162 
163<!-- npu="950" id8 -->163<!-- npu="950" id8 -->
164- 
165以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:164以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
166 165 
167```bash166```bash
@@ -86,7 +86,6 @@ __simd_callee__ inline void asc_loadalign_brc_datablock_postupdate(vector_int8_t
86将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。86将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
87 87 
88<!-- npu="950" id8 -->88<!-- npu="950" id8 -->
89- 
90以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:89以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
91 90 
92```bash91```bash
@@ -161,7 +161,6 @@ __simd_callee__ inline void asc_loadalign_brc_elem(vector_half& dst,
161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
162 162 
163<!-- npu="950" id8 -->163<!-- npu="950" id8 -->
164- 
165以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:164以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
166 165 
167```bash166```bash
@@ -161,7 +161,6 @@ __simd_callee__ inline void asc_loadalign_brc_elem2datablock(vector_half& dst,
161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
162 162 
163<!-- npu="950" id8 -->163<!-- npu="950" id8 -->
164- 
165以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:164以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
166 165 
167```bash166```bash
@@ -84,7 +84,6 @@ __simd_callee__ inline void asc_loadalign_brc_elem2datablock_postupdate(vector_h
84将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。84将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
85 85 
86<!-- npu="950" id8 -->86<!-- npu="950" id8 -->
87- 
88以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:87以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
89 88 
90```bash89```bash
@@ -86,7 +86,6 @@ __simd_callee__ inline void asc_loadalign_brc_elem_postupdate(vector_int16_t& ds
86将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。86将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
87 87 
88<!-- npu="950" id8 -->88<!-- npu="950" id8 -->
89- 
90以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:89以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
91 90 
92```bash91```bash
@@ -88,7 +88,6 @@ __simd_callee__ inline vector_uint8_t asc_loadalign_datablock_strided(__ubuf__ u
88将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。88将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
89 89 
90<!-- npu="950" id8 -->90<!-- npu="950" id8 -->
91- 
92以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:91以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
93 92 
94```bash93```bash
@@ -165,7 +165,6 @@ __simd_callee__ inline void asc_loadalign_deintlv(vector_float& dst0,
165将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。165将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
166 166 
167<!-- npu="950" id8 -->167<!-- npu="950" id8 -->
168- 
169以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:168以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
170 169 
171```bash170```bash
@@ -91,7 +91,6 @@ __simd_callee__ inline void asc_loadalign_deintlv_postupdate(vector_int16_t& dst
91将代码保存为`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)。
92 92 
93<!-- npu="950" id8 -->93<!-- npu="950" id8 -->
94- 
95以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:94以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
96 95 
97```bash96```bash
@@ -181,7 +181,6 @@ __simd_callee__ inline void asc_loadalign_downsample(vector_half& dst,
181将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。181将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
182 182 
183<!-- npu="950" id8 -->183<!-- npu="950" id8 -->
184- 
185以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:184以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
186 185 
187```bash186```bash
@@ -92,7 +92,6 @@ __simd_callee__ inline void asc_loadalign_downsample_postupdate(vector_int16_t&
92将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。92将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
93 93 
94<!-- npu="950" id8 -->94<!-- npu="950" id8 -->
95- 
96以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:95以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
97 96 
98```bash97```bash
@@ -64,7 +64,6 @@ __simd_callee__ inline vector_bool asc_loadalign_mask(__ubuf__ uint32_t* src)
64将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。64将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
65 65 
66<!-- npu="950" id8 -->66<!-- npu="950" id8 -->
67- 
68以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:67以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
69 68 
70```bash69```bash
@@ -64,7 +64,6 @@ __simd_callee__ inline vector_bool asc_loadalign_mask_downsample(__ubuf__ uint32
64将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。64将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
65 65 
66<!-- npu="950" id8 -->66<!-- npu="950" id8 -->
67- 
68以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:67以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
69 68 
70```bash69```bash
@@ -64,7 +64,6 @@ __simd_callee__ inline vector_bool asc_loadalign_mask_upsample(__ubuf__ uint32_t
64将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。64将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
65 65 
66<!-- npu="950" id8 -->66<!-- npu="950" id8 -->
67- 
68以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:67以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
69 68 
70```bash69```bash
@@ -161,7 +161,6 @@ __simd_callee__ inline void asc_loadalign_unpack(vector_half& dst,
161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
162 162 
163<!-- npu="950" id8 -->163<!-- npu="950" id8 -->
164- 
165以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:164以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
166 165 
167```bash166```bash
@@ -161,7 +161,6 @@ __simd_callee__ inline void asc_loadalign_unpack4(vector_int8_t& dst,
161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。161将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
162 162 
163<!-- npu="950" id8 -->163<!-- npu="950" id8 -->
164- 
165以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:164以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
166 165 
167```bash166```bash
@@ -84,7 +84,6 @@ __simd_callee__ inline void asc_loadalign_unpack4_postupdate(vector_int8_t& dst,
84将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。84将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
85 85 
86<!-- npu="950" id8 -->86<!-- npu="950" id8 -->
87- 
88以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:87以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
89 88 
90```bash89```bash
@@ -84,7 +84,6 @@ __simd_callee__ inline void asc_loadalign_unpack_postupdate(vector_int16_t& dst,
84将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。84将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
85 85 
86<!-- npu="950" id8 -->86<!-- npu="950" id8 -->
87- 
88以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:87以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
89 88 
90```bash89```bash
@@ -181,7 +181,6 @@ __simd_callee__ inline void asc_loadalign_upsample(vector_half& dst,
181将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。181将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
182 182 
183<!-- npu="950" id8 -->183<!-- npu="950" id8 -->
184- 
185以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:184以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
186 185 
187```bash186```bash
@@ -92,7 +92,6 @@ __simd_callee__ inline void asc_loadalign_upsample_postupdate(vector_half& dst,
92将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。92将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
93 93 
94<!-- npu="950" id8 -->94<!-- npu="950" id8 -->
95- 
96以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:95以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
97 96 
98```bash97```bash
@@ -81,7 +81,6 @@ __simd_callee__ inline void asc_loadunalign(vector_int8_t& dst,
81将代码保存为`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)。
82 82 
83<!-- npu="950" id8 -->83<!-- npu="950" id8 -->
84- 
85以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:84以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
86 85 
87```bash86```bash
@@ -159,7 +159,6 @@ __simd_callee__ inline void asc_loadunalign_postupdate(vector_int8_t& dst,
159将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。159将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
160 160 
161<!-- npu="950" id8 -->161<!-- npu="950" id8 -->
162- 
163以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:162以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
164 163 
165```bash164```bash
@@ -149,7 +149,6 @@ __simd_callee__ inline void asc_loadunalign_pre(vector_load_unalign& dst,
149将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。149将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
150 150 
151<!-- npu="950" id8 -->151<!-- npu="950" id8 -->
152- 
153以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:152以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
154 153 
155```bash154```bash
@@ -43,7 +43,7 @@
43 43 
44- 1维偏移44- 1维偏移
45 45 
46- ```cpp46+ ```c
47 __simd_callee__ inline addr_reg asc_update_addr_reg_b8(uint32_t offset)47 __simd_callee__ inline addr_reg asc_update_addr_reg_b8(uint32_t offset)
48 __simd_callee__ inline addr_reg asc_update_addr_reg_b16(uint32_t offset)48 __simd_callee__ inline addr_reg asc_update_addr_reg_b16(uint32_t offset)
49 __simd_callee__ inline addr_reg asc_update_addr_reg_b32(uint32_t offset)49 __simd_callee__ inline addr_reg asc_update_addr_reg_b32(uint32_t offset)
@@ -51,26 +51,44 @@
51 51 
52- 2维偏移52- 2维偏移
53 53 
54- ```cpp54+ ```c
55- __simd_callee__ inline addr_reg asc_update_addr_reg_b8(uint32_t offset0, uint32_t offset1)55+ __simd_callee__ inline addr_reg asc_update_addr_reg_b8(uint32_t offset0,
56- __simd_callee__ inline addr_reg asc_update_addr_reg_b16(uint32_t offset0, uint32_t offset1)56+ uint32_t offset1)
57- __simd_callee__ inline addr_reg asc_update_addr_reg_b32(uint32_t offset0, uint32_t offset1)57+ __simd_callee__ inline addr_reg asc_update_addr_reg_b16(uint32_t offset0,
58+ uint32_t offset1)
59+ __simd_callee__ inline addr_reg asc_update_addr_reg_b32(uint32_t offset0,
60+ uint32_t offset1)
58 ```61 ```
59 62 
60- 3维偏移63- 3维偏移
61 64 
62- ```cpp65+ ```c
63- __simd_callee__ inline addr_reg asc_update_addr_reg_b8(uint32_t offset0, uint32_t offset1, uint32_t offset2)66+ __simd_callee__ inline addr_reg asc_update_addr_reg_b8(uint32_t offset0,
64- __simd_callee__ inline addr_reg asc_update_addr_reg_b16(uint32_t offset0, uint32_t offset1, uint32_t offset2)67+ uint32_t offset1,
65- __simd_callee__ inline addr_reg asc_update_addr_reg_b32(uint32_t offset0, uint32_t offset1, uint32_t offset2)68+ uint32_t offset2)
69+ __simd_callee__ inline addr_reg asc_update_addr_reg_b16(uint32_t offset0,
70+ uint32_t offset1,
71+ uint32_t offset2)
72+ __simd_callee__ inline addr_reg asc_update_addr_reg_b32(uint32_t offset0,
73+ uint32_t offset1,
74+ uint32_t offset2)
66 ```75 ```
67 76 
68- 4维偏移77- 4维偏移
69 78 
70- ```cpp79+ ```c
71- __simd_callee__ inline addr_reg asc_update_addr_reg_b8(uint32_t offset0, uint32_t offset1, uint32_t offset2, uint32_t offset3)80+ __simd_callee__ inline addr_reg asc_update_addr_reg_b8(uint32_t offset0,
72- __simd_callee__ inline addr_reg asc_update_addr_reg_b16(uint32_t offset0, uint32_t offset1, uint32_t offset2, uint32_t offset3)81+ uint32_t offset1,
73- __simd_callee__ inline addr_reg asc_update_addr_reg_b32(uint32_t offset0, uint32_t offset1, uint32_t offset2, uint32_t offset3)82+ uint32_t offset2,
83+ uint32_t offset3)
84+ __simd_callee__ inline addr_reg asc_update_addr_reg_b16(uint32_t offset0,
85+ uint32_t offset1,
86+ uint32_t offset2,
87+ uint32_t offset3)
88+ __simd_callee__ inline addr_reg asc_update_addr_reg_b32(uint32_t offset0,
89+ uint32_t offset1,
90+ uint32_t offset2,
91+ uint32_t offset3)
74 ```92 ```
75 93 
76## 参数说明94## 参数说明
@@ -99,16 +117,96 @@
99 117 
100## 调用示例118## 调用示例
101 119 
120+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
121+ 
122+<!-- npu="950" id8 -->
M
Mmunanhw10 天前

结束标签位置不准确

likedislike
123+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
124+ 
125+```bash
126+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
127+```
128+ 
129+<!-- end id8 -->
130+ 
102```cpp131```cpp
103-__simd_vf__ inline void update_addr_reg_vf(__ubuf__ half* dst_addr, __ubuf__ half* src_addr, uint16_t one_repeat_size, uint16_t repeat_time)132+#include <cstdint>
133+#include <iostream>
134+#include <vector>
135+#include "c_api/asc_simd.h"
136+#include "acl/acl.h"
137+ 
138+namespace {
139+template <typename T>
140+bool compare_data(const std::vector<T>& actual, const std::vector<T>& expected, double tolerance = 0.0)
104{141{
105- vector_half data;142+ if (actual.size() != expected.size()) return false;
106- vector_bool mask = asc_create_mask_b16(PAT_ALL);143+ for (size_t i = 0; i < actual.size(); ++i) {
107- addr_reg offset;144+ if (actual[i] == expected[i]) continue;
108- for (uint16_t i = 0; i < repeat_time; ++i) {145+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
109- offset = asc_update_addr_reg_b16(one_repeat_size);146+ if (diff > tolerance || diff < -tolerance) return false;
110- asc_loadalign(data, src_addr, offset);147+ }
111- asc_storealign(dst_addr, data, offset, mask);148+ return true;
149+}
150+ 
151+constexpr uint16_t ONE_REPEAT_SIZE = 256;
152+constexpr uint16_t REPEAT_COUNT = 4;
153+constexpr uint32_t TOTAL_BYTES = ONE_REPEAT_SIZE * REPEAT_COUNT;
154+ 
155+__simd_vf__ inline void exercise(__ubuf__ uint8_t* output, __ubuf__ uint8_t* input,
156+ uint16_t oneRepeatSize, uint16_t repeatTimes)
157+{
158+ vector_bool mask = asc_create_mask_b8(PAT_ALL);
159+ vector_uint8_t value;
160+ addr_reg offset_reg;
161+ for (uint16_t i = 0; i < repeatTimes; ++i) {
162+ offset_reg = asc_update_addr_reg_b8(oneRepeatSize);
163+ asc_loadalign(value, input, offset_reg);
164+ asc_storealign(output, value, offset_reg, mask);
112 }165 }
113}166}
167+ 
168+__global__ __vector__ void asc_update_addr_reg_kernel(__gm__ uint8_t* output, __gm__ uint8_t* input,
169+ uint16_t oneRepeatSize, uint16_t repeatTimes)
170+{
171+ asc_init();
172+ __ubuf__ uint8_t output_local[TOTAL_BYTES];
173+ __ubuf__ uint8_t input_local[TOTAL_BYTES];
174+ asc_copy_gm2ub_align(input_local, input, TOTAL_BYTES);
175+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
176+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
177+ exercise(output_local, input_local, oneRepeatSize, repeatTimes);
178+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
179+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
180+ asc_copy_ub2gm_align(output, output_local, TOTAL_BYTES);
181+ asc_sync();
182+}
183+} // namespace
184+ 
185+int main()
186+{
187+ std::vector<uint8_t> input(TOTAL_BYTES), output(TOTAL_BYTES, 0xff), golden(TOTAL_BYTES);
188+ for (uint32_t i = 0; i < TOTAL_BYTES; ++i) input[i] = static_cast<uint8_t>(i % 251);
189+ golden = input;
190+ aclInit(nullptr);
191+ aclrtSetDevice(0);
192+ uint8_t* input_device = nullptr;
193+ aclrtMalloc(reinterpret_cast<void**>(&input_device), TOTAL_BYTES * sizeof(uint8_t),
194+ ACL_MEM_MALLOC_HUGE_FIRST);
195+ uint8_t* output_device = nullptr;
196+ aclrtMalloc(reinterpret_cast<void**>(&output_device), TOTAL_BYTES * sizeof(uint8_t),
197+ ACL_MEM_MALLOC_HUGE_FIRST);
198+ aclrtMemcpy(input_device, input.size() * sizeof(uint8_t), input.data(), input.size() * sizeof(uint8_t),
199+ ACL_MEMCPY_HOST_TO_DEVICE);
200+ asc_update_addr_reg_kernel<<<1, 0>>>(output_device, input_device, ONE_REPEAT_SIZE, REPEAT_COUNT);
201+ aclrtSynchronizeDevice();
202+ aclrtMemcpy(output.data(), output.size() * sizeof(uint8_t), output_device, output.size() * sizeof(uint8_t),
203+ ACL_MEMCPY_DEVICE_TO_HOST);
204+ const bool passed = compare_data(output, golden);
205+ std::cout << (passed ? "[Success] asc_update_addr_reg_b8 passed." : "[Failed] asc_update_addr_reg_b8 failed.") << std::endl;
206+ aclrtFree(input_device);
207+ aclrtFree(output_device);
208+ aclrtResetDevice(0);
209+ aclFinalize();
210+ return passed ? 0 : 1;
211+}
114```212```
@@ -88,14 +88,15 @@ __simd_callee__ inline void asc_abs(vector_half& dst,
88 88 
89## 调用示例89## 调用示例
90 90 
91-<!-- npu="950" id8 -->
92将以下代码保存为`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)。
93 92 
93+<!-- npu="950" id8 -->
94以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:94以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
95 95 
96```bash96```bash
97bisheng example.asc -o main --npu-arch=dav-3510 && ./main97bisheng example.asc -o main --npu-arch=dav-3510 && ./main
98```98```
99+<!-- end id8 -->
99 100 
100```cpp101```cpp
101#include <algorithm>102#include <algorithm>
@@ -182,5 +183,3 @@ int main()
182 return passed ? 0 : 1;183 return passed ? 0 : 1;
183}184}
184```185```
185- 
186-<!-- end id8 -->
@@ -91,14 +91,15 @@ __simd_callee__ inline void asc_add_scalar(vector_half& dst,
91 91 
92## 调用示例92## 调用示例
93 93 
94-<!-- npu="950" id8 -->
95将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。94将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
96 95 
96+<!-- npu="950" id8 -->
97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
98 98 
99```bash99```bash
100bisheng example.asc -o main --npu-arch=dav-3510 && ./main100bisheng example.asc -o main --npu-arch=dav-3510 && ./main
101```101```
102+<!-- end id8 -->
102 103 
103```cpp104```cpp
104#include <algorithm>105#include <algorithm>
@@ -189,5 +190,3 @@ int main()
189 return passed ? 0 : 1;190 return passed ? 0 : 1;
190}191}
191```192```
192- 
193-<!-- end id8 -->
@@ -78,14 +78,15 @@ __simd_callee__ inline void asc_ln(vector_half& dst,
78 78 
79## 调用示例79## 调用示例
80 80 
81-<!-- npu="950" id8 -->
82将以下代码保存为`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)。
83 82 
83+<!-- npu="950" id8 -->
84以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:84以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
85 85 
86```bash86```bash
87bisheng example.asc -o main --npu-arch=dav-3510 -lm && ./main87bisheng example.asc -o main --npu-arch=dav-3510 -lm && ./main
88```88```
89+<!-- end id8 -->
89 90 
90```cpp91```cpp
91#include <algorithm>92#include <algorithm>
@@ -175,5 +176,3 @@ int main()
175 return passed ? 0 : 1;176 return passed ? 0 : 1;
176}177}
177```178```
178- 
179-<!-- end id8 -->
@@ -77,16 +77,17 @@ __simd_callee__ inline void asc_log(vector_half& dst,
77- `dst``src`的数据类型需要保持一致。77- `dst``src`的数据类型需要保持一致。
78- `mask`掩码位为0时,`dst`对应元素置0。78- `mask`掩码位为0时,`dst`对应元素置0。
79 79 
80-<!-- npu="950" id8 -->
81## 调用示例80## 调用示例
82 81 
83将以下代码保存为`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)。
84 83 
84+<!-- npu="950" id8 -->
85以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:85以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
86 86 
87```bash87```bash
88bisheng example.asc -o main --npu-arch=dav-3510 -lm && ./main88bisheng example.asc -o main --npu-arch=dav-3510 -lm && ./main
89```89```
90+<!-- end id8 -->
90 91 
91```cpp92```cpp
92#include <algorithm>93#include <algorithm>
@@ -176,5 +177,3 @@ int main()
176 return passed ? 0 : 1;177 return passed ? 0 : 1;
177}178}
178```179```
179- 
180-<!-- end id8 -->
@@ -92,14 +92,15 @@ __simd_callee__ inline void asc_max_scalar(vector_half& dst,
92 92 
93## 调用示例93## 调用示例
94 94 
95-<!-- npu="950" id8 -->
96将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。95将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
97 96 
97+<!-- npu="950" id8 -->
98以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:98以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
99 99 
100```bash100```bash
101bisheng example.asc -o main --npu-arch=dav-3510 && ./main101bisheng example.asc -o main --npu-arch=dav-3510 && ./main
102```102```
103+<!-- end id8 -->
103 104 
104```cpp105```cpp
105#include <algorithm>106#include <algorithm>
@@ -190,5 +191,3 @@ int main()
190 return passed ? 0 : 1;191 return passed ? 0 : 1;
191}192}
192```193```
193- 
194-<!-- end id8 -->
@@ -91,14 +91,15 @@ __simd_callee__ inline void asc_min_scalar(vector_half& dst,
91 91 
92## 调用示例92## 调用示例
93 93 
94-<!-- npu="950" id8 -->
95将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。94将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
96 95 
96+<!-- npu="950" id8 -->
97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
98 98 
99```bash99```bash
100bisheng example.asc -o main --npu-arch=dav-3510 && ./main100bisheng example.asc -o main --npu-arch=dav-3510 && ./main
101```101```
102+<!-- end id8 -->
102 103 
103```cpp104```cpp
104#include <algorithm>105#include <algorithm>
@@ -189,5 +190,3 @@ int main()
189 return passed ? 0 : 1;190 return passed ? 0 : 1;
190}191}
191```192```
192- 
193-<!-- end id8 -->
@@ -89,14 +89,15 @@ __simd_callee__ inline void asc_mul_scalar(vector_half& dst,
89 89 
90## 调用示例90## 调用示例
91 91 
92-<!-- npu="950" id8 -->
93将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。92将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
94 93 
94+<!-- npu="950" id8 -->
95以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:95以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
96 96 
97```bash97```bash
98bisheng example.asc -o main --npu-arch=dav-3510 && ./main98bisheng example.asc -o main --npu-arch=dav-3510 && ./main
99```99```
100+<!-- end id8 -->
100 101 
101```cpp102```cpp
102#include <algorithm>103#include <algorithm>
@@ -187,5 +188,3 @@ int main()
187 return passed ? 0 : 1;188 return passed ? 0 : 1;
188}189}
189```190```
190- 
191-<!-- end id8 -->
@@ -84,16 +84,17 @@ __simd_callee__ inline void asc_duplicate(vector_int8_t& dst,
84- 使用mask前,需要通过掩码设置或搬入接口完成初始化;未初始化的掩码寄存器内容不确定。84- 使用mask前,需要通过掩码设置或搬入接口完成初始化;未初始化的掩码寄存器内容不确定。
85- mask仅筛选dst中写入广播值的位置,不筛选src中的元素。无论mask的最低位是否有效,待广播的数据均为src的最低位元素。85- mask仅筛选dst中写入广播值的位置,不筛选src中的元素。无论mask的最低位是否有效,待广播的数据均为src的最低位元素。
86 86 
87-<!-- npu="950" id8 -->
88## 调用示例87## 调用示例
89 88 
90将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。89将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
91 90 
91+<!-- npu="950" id8 -->
92以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:92以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
93 93 
94```bash94```bash
95bisheng example.asc -o main --npu-arch=dav-3510 && ./main95bisheng example.asc -o main --npu-arch=dav-3510 && ./main
96```96```
97+<!-- end id8 -->
97 98 
98以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。99以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
99 100 
@@ -167,4 +168,3 @@ int main()
167 return passed ? 0 : 1;168 return passed ? 0 : 1;
168}169}
169```170```
170-<!-- end id8 -->
@@ -115,16 +115,17 @@ __simd_callee__ inline void asc_duplicate_scalar(vector_int8_t& dst,
115- 使用mask前,需要通过掩码设置或搬入接口完成初始化;未初始化的掩码寄存器内容不确定。115- 使用mask前,需要通过掩码设置或搬入接口完成初始化;未初始化的掩码寄存器内容不确定。
116- 调用带返回值接口时,应使用类型明确的变量或显式类型转换,例如`static_cast<uint8_t>(1)`,以匹配正确的函数原型。116- 调用带返回值接口时,应使用类型明确的变量或显式类型转换,例如`static_cast<uint8_t>(1)`,以匹配正确的函数原型。
117 117 
118-<!-- npu="950" id8 -->
119## 调用示例118## 调用示例
120 119 
121将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。120将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
122 121 
122+<!-- npu="950" id8 -->
123以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:123以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
124 124 
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+<!-- end id8 -->
128 129 
129以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。130以下调用示例代码仅Ascend 950PR/Ascend 950DT产品支持。
130 131 
@@ -197,4 +198,3 @@ int main()
197 return passed ? 0 : 1;198 return passed ? 0 : 1;
198}199}
199```200```
200-<!-- end id8 -->
@@ -96,14 +96,15 @@ __simd_callee__ inline void asc_ge_scalar(vector_bool& dst,
96 96 
97## 调用示例97## 调用示例
98 98 
99-<!-- npu="950" id8 -->
100将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。99将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
101 100 
101+<!-- npu="950" id8 -->
102以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:102以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
103 103 
104```bash104```bash
105bisheng example.asc -o main --npu-arch=dav-3510 && ./main105bisheng example.asc -o main --npu-arch=dav-3510 && ./main
106```106```
107+<!-- end id8 -->
107 108 
108```cpp109```cpp
109#include <cstdint>110#include <cstdint>
@@ -194,5 +195,3 @@ int main()
194 return passed ? 0 : 1;195 return passed ? 0 : 1;
195}196}
196```197```
197- 
198-<!-- end id8 -->
@@ -94,14 +94,15 @@ __simd_callee__ inline void asc_gt_scalar(vector_bool& dst,
94 94 
95## 调用示例95## 调用示例
96 96 
97-<!-- npu="950" id8 -->
98将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。97将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
99 98 
99+<!-- npu="950" id8 -->
100以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:100以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
101 101 
102```bash102```bash
103bisheng example.asc -o main --npu-arch=dav-3510 && ./main103bisheng example.asc -o main --npu-arch=dav-3510 && ./main
104```104```
105+<!-- end id8 -->
105 106 
106```cpp107```cpp
107#include <cstdint>108#include <cstdint>
@@ -192,5 +193,3 @@ int main()
192 return passed ? 0 : 1;193 return passed ? 0 : 1;
193}194}
194```195```
195- 
196-<!-- end id8 -->
@@ -96,14 +96,15 @@ __simd_callee__ inline void asc_ne_scalar(vector_bool& dst,
96 96 
97## 调用示例97## 调用示例
98 98 
99-<!-- npu="950" id8 -->
100将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。99将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
101 100 
101+<!-- npu="950" id8 -->
102以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:102以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
103 103 
104```bash104```bash
105bisheng example.asc -o main --npu-arch=dav-3510 && ./main105bisheng example.asc -o main --npu-arch=dav-3510 && ./main
106```106```
107+<!-- end id8 -->
107 108 
108```cpp109```cpp
109#include <cstdint>110#include <cstdint>
@@ -194,5 +195,3 @@ int main()
194 return passed ? 0 : 1;195 return passed ? 0 : 1;
195}196}
196```197```
197- 
198-<!-- end id8 -->
@@ -103,16 +103,15 @@ __simd_callee__ inline void asc_float2bfloat16_rn(vector_bfloat16_t& dst,
103 103 
104## 调用示例104## 调用示例
105 105 
106- 
107-<!-- npu="950" id8 -->
108- 
109将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。106将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
110 107 
108+<!-- npu="950" id8 -->
111以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:109以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
112 110 
113```bash111```bash
114bisheng example.asc -o main --npu-arch=dav-3510 && ./main112bisheng example.asc -o main --npu-arch=dav-3510 && ./main
115```113```
114+<!-- end id8 -->
116 115 
117```c116```c
118#include <cstdint>117#include <cstdint>
@@ -334,5 +333,3 @@ int main()
334 return passed ? 0 : 1;333 return passed ? 0 : 1;
335}334}
336```335```
337- 
338-<!-- end id8 -->
@@ -102,16 +102,15 @@ __simd_callee__ inline void asc_float2half_rn(vector_half& dst,
102 102 
103## 调用示例103## 调用示例
104 104 
105+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
105 106 
106<!-- npu="950" id8 -->107<!-- npu="950" id8 -->
107- 
108-f代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
109- 
110以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:108以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
111 109 
112```bash110```bash
113bisheng example.asc -o main --npu-arch=dav-3510 && ./main111bisheng example.asc -o main --npu-arch=dav-3510 && ./main
114```112```
113+<!-- end id8 -->
115 114 
116```c115```c
117#include <cstdint>116#include <cstdint>
@@ -212,5 +211,3 @@ int main()
212 return passed ? 0 : 1;211 return passed ? 0 : 1;
213}212}
214```213```
215- 
216-<!-- end id8 -->
@@ -104,16 +104,15 @@ __simd_callee__ inline void asc_float2int16_rn(vector_int16_t& dst,
104 104 
105## 调用示例105## 调用示例
106 106 
107+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
107 108 
108<!-- npu="950" id8 -->109<!-- npu="950" id8 -->
109- 
110-f代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
111- 
112以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:110以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
113 111 
114```bash112```bash
115bisheng example.asc -o main --npu-arch=dav-3510 && ./main113bisheng example.asc -o main --npu-arch=dav-3510 && ./main
116```114```
115+<!-- end id8 -->
117 116 
118```c117```c
119#include <cstdint>118#include <cstdint>
@@ -214,5 +213,3 @@ int main()
214 return passed ? 0 : 1;213 return passed ? 0 : 1;
215}214}
216```215```
217- 
218-<!-- end id8 -->
@@ -93,16 +93,15 @@ __simd_callee__ inline void asc_float2int32_rn(vector_int32_t& dst,
93 93 
94## 调用示例94## 调用示例
95 95 
96+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
96 97 
97<!-- npu="950" id8 -->98<!-- npu="950" id8 -->
98- 
99-f代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
100- 
101以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:99以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
102 100 
103```bash101```bash
104bisheng example.asc -o main --npu-arch=dav-3510 && ./main102bisheng example.asc -o main --npu-arch=dav-3510 && ./main
105```103```
104+<!-- end id8 -->
106 105 
107```c106```c
108#include <cstdint>107#include <cstdint>
@@ -201,5 +200,3 @@ int main()
201 return passed ? 0 : 1;200 return passed ? 0 : 1;
202}201}
203```202```
204- 
205-<!-- end id8 -->
@@ -103,16 +103,15 @@ __simd_callee__ inline void asc_float2int64_rn(vector_int64_t& dst,
103 103 
104## 调用示例104## 调用示例
105 105 
106+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
106 107 
107<!-- npu="950" id8 -->108<!-- npu="950" id8 -->
108- 
109-f代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
110- 
111以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:109以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
112 110 
113```bash111```bash
114bisheng example.asc -o main --npu-arch=dav-3510 && ./main112bisheng example.asc -o main --npu-arch=dav-3510 && ./main
115```113```
114+<!-- end id8 -->
116 115 
117```c116```c
118#include <cstdint>117#include <cstdint>
@@ -213,5 +212,3 @@ int main()
213 return passed ? 0 : 1;212 return passed ? 0 : 1;
214}213}
215```214```
216- 
217-<!-- end id8 -->
@@ -85,16 +85,15 @@ __simd_callee__ inline void asc_hif82half(vector_half& dst,
85 85 
86## 调用示例86## 调用示例
87 87 
88+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
88 89 
89<!-- npu="950" id8 -->90<!-- npu="950" id8 -->
90- 
91-s代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
92- 
93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:91以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
94 92 
95```bash93```bash
96bisheng example.asc -o main --npu-arch=dav-3510 && ./main94bisheng example.asc -o main --npu-arch=dav-3510 && ./main
97```95```
96+<!-- end id8 -->
98 97 
99```c98```c
100#include <cstdint>99#include <cstdint>
@@ -190,5 +189,3 @@ int main()
190 return passed ? 0 : 1;189 return passed ? 0 : 1;
191}190}
192```191```
193- 
194-<!-- end id8 -->
@@ -85,16 +85,15 @@ __simd_callee__ inline void asc_int162float(vector_float& dst,
85 85 
86## 调用示例86## 调用示例
87 87 
88+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
88 89 
89<!-- npu="950" id8 -->90<!-- npu="950" id8 -->
90- 
91-s代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
92- 
93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:91以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
94 92 
95```bash93```bash
96bisheng example.asc -o main --npu-arch=dav-3510 && ./main94bisheng example.asc -o main --npu-arch=dav-3510 && ./main
97```95```
96+<!-- end id8 -->
98 97 
99```c98```c
100#include <cstdint>99#include <cstdint>
@@ -190,5 +189,3 @@ int main()
190 return passed ? 0 : 1;189 return passed ? 0 : 1;
191}190}
192```191```
193- 
194-<!-- end id8 -->
@@ -95,16 +95,15 @@ __simd_callee__ inline void asc_int162half_rn(vector_half& dst,
95 95 
96## 调用示例96## 调用示例
97 97 
98+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
98 99 
99<!-- npu="950" id8 -->100<!-- npu="950" id8 -->
100- 
101-s代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
102- 
103以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:101以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
104 102 
105```bash103```bash
106bisheng example.asc -o main --npu-arch=dav-3510 && ./main104bisheng example.asc -o main --npu-arch=dav-3510 && ./main
107```105```
106+<!-- end id8 -->
108 107 
109```c108```c
110#include <cstdint>109#include <cstdint>
@@ -200,5 +199,3 @@ int main()
200 return passed ? 0 : 1;199 return passed ? 0 : 1;
201}200}
202```201```
203- 
204-<!-- end id8 -->
@@ -85,16 +85,15 @@ __simd_callee__ inline void asc_int162int32(vector_int32_t& dst,
85 85 
86## 调用示例86## 调用示例
87 87 
88+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
88 89 
89<!-- npu="950" id8 -->90<!-- npu="950" id8 -->
90- 
91-s代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
92- 
93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:91以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
94 92 
95```bash93```bash
96bisheng example.asc -o main --npu-arch=dav-3510 && ./main94bisheng example.asc -o main --npu-arch=dav-3510 && ./main
97```95```
96+<!-- end id8 -->
98 97 
99```c98```c
100#include <cstdint>99#include <cstdint>
@@ -190,5 +189,3 @@ int main()
190 return passed ? 0 : 1;189 return passed ? 0 : 1;
191}190}
192```191```
193- 
194-<!-- end id8 -->
@@ -85,16 +85,15 @@ __simd_callee__ inline void asc_int162uint32(vector_uint32_t& dst,
85 85 
86## 调用示例86## 调用示例
87 87 
88+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
88 89 
89<!-- npu="950" id8 -->90<!-- npu="950" id8 -->
90- 
91-s代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
92- 
93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:91以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
94 92 
95```bash93```bash
96bisheng example.asc -o main --npu-arch=dav-3510 && ./main94bisheng example.asc -o main --npu-arch=dav-3510 && ./main
97```95```
96+<!-- end id8 -->
98 97 
99```c98```c
100#include <cstdint>99#include <cstdint>
@@ -192,5 +191,3 @@ int main()
192 return passed ? 0 : 1;191 return passed ? 0 : 1;
193}192}
194```193```
195- 
196-<!-- end id8 -->
@@ -101,16 +101,15 @@ __simd_callee__ inline void asc_int162uint8(vector_uint8_t& dst,
101 101 
102## 调用示例102## 调用示例
103 103 
104+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
104 105 
105<!-- npu="950" id8 -->106<!-- npu="950" id8 -->
106- 
107-s代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
108- 
109以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:107以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
110 108 
111```bash109```bash
112bisheng example.asc -o main --npu-arch=dav-3510 && ./main110bisheng example.asc -o main --npu-arch=dav-3510 && ./main
113```111```
112+<!-- end id8 -->
114 113 
115```c114```c
116#include <cstdint>115#include <cstdint>
@@ -214,5 +213,3 @@ int main()
214 return passed ? 0 : 1;213 return passed ? 0 : 1;
215}214}
216```215```
217- 
218-<!-- end id8 -->
@@ -85,15 +85,15 @@ __simd_callee__ inline void asc_int322float_rn(vector_float& dst,
85 85 
86## 调用示例86## 调用示例
87 87 
88-<!-- npu="950" id8 -->
89- 
90将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。88将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
91 89 
90+<!-- npu="950" id8 -->
92以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:91以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
93 92 
94```bash93```bash
95bisheng example.asc -o main --npu-arch=dav-3510 && ./main94bisheng example.asc -o main --npu-arch=dav-3510 && ./main
96```95```
96+<!-- end id8 -->
97 97 
98```c98```c
99#include <cstdint>99#include <cstdint>
@@ -187,4 +187,3 @@ int main()
187}187}
188```188```
189 189 
190-<!-- end id8 -->
@@ -96,15 +96,15 @@ __simd_callee__ inline void asc_int322int16(vector_int16_t& dst,
96 96 
97## 调用示例97## 调用示例
98 98 
99-<!-- npu="950" id8 -->
100- 
101将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。99将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
102 100 
101+<!-- npu="950" id8 -->
103以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:102以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
104 103 
105```bash104```bash
106bisheng example.asc -o main --npu-arch=dav-3510 && ./main105bisheng example.asc -o main --npu-arch=dav-3510 && ./main
107```106```
107+<!-- end id8 -->
108 108 
109```c109```c
110#include <cstdint>110#include <cstdint>
@@ -223,4 +223,3 @@ int main()
223}223}
224```224```
225 225 
226-<!-- end id8 -->
@@ -78,16 +78,15 @@ __simd_callee__ inline void asc_int322int64(vector_int64_t& dst,
78 78 
79## 调用示例79## 调用示例
80 80 
81- 
82-<!-- npu="950" id8 -->
83- 
84将代码保存为`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)。
85 82 
83+<!-- npu="950" id8 -->
86以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:84以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
87 85 
88```bash86```bash
89bisheng example.asc -o main --npu-arch=dav-3510 && ./main87bisheng example.asc -o main --npu-arch=dav-3510 && ./main
90```88```
89+<!-- end id8 -->
91 90 
92```c91```c
93#include <cstdint>92#include <cstdint>
@@ -191,5 +190,3 @@ int main()
191 return passed ? 0 : 1;190 return passed ? 0 : 1;
192}191}
193```192```
194- 
195-<!-- end id8 -->
@@ -96,16 +96,15 @@ __simd_callee__ inline void asc_int322uint16(vector_uint16_t& dst,
96 96 
97## 调用示例97## 调用示例
98 98 
99- 
100-<!-- npu="950" id8 -->
101- 
102将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。99将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
103 100 
101+<!-- npu="950" id8 -->
104以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:102以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
105 103 
106```bash104```bash
107bisheng example.asc -o main --npu-arch=dav-3510 && ./main105bisheng example.asc -o main --npu-arch=dav-3510 && ./main
108```106```
107+<!-- end id8 -->
109 108 
110```c109```c
111#include <cstdint>110#include <cstdint>
@@ -223,5 +222,3 @@ int main()
223 return passed ? 0 : 1;222 return passed ? 0 : 1;
224}223}
225```224```
226- 
227-<!-- end id8 -->
@@ -104,16 +104,15 @@ __simd_callee__ inline void asc_int642float_rn(vector_float& dst,
104 104 
105## 调用示例105## 调用示例
106 106 
107- 
108-<!-- npu="950" id8 -->
109- 
110将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。107将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
111 108 
109+<!-- npu="950" id8 -->
112以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:110以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
113 111 
114```bash112```bash
115bisheng example.asc -o main --npu-arch=dav-3510 && ./main113bisheng example.asc -o main --npu-arch=dav-3510 && ./main
116```114```
115+<!-- end id8 -->
117 116 
118```c117```c
119#include <cstdint>118#include <cstdint>
@@ -209,5 +208,3 @@ int main()
209 return passed ? 0 : 1;208 return passed ? 0 : 1;
210}209}
211```210```
212- 
213-<!-- end id8 -->
@@ -103,16 +103,15 @@ __simd_callee__ inline void asc_int642int32(vector_int32_t& dst,
103 103 
104## 调用示例104## 调用示例
105 105 
106- 
107-<!-- npu="950" id8 -->
108- 
109将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。106将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
110 107 
108+<!-- npu="950" id8 -->
111以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:109以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
112 110 
113```bash111```bash
114bisheng example.asc -o main --npu-arch=dav-3510 && ./main112bisheng example.asc -o main --npu-arch=dav-3510 && ./main
115```113```
114+<!-- end id8 -->
116 115 
117```c116```c
118#include <cstdint>117#include <cstdint>
@@ -208,5 +207,3 @@ int main()
208 return passed ? 0 : 1;207 return passed ? 0 : 1;
209}208}
210```209```
211- 
212-<!-- end id8 -->
@@ -85,16 +85,15 @@ __simd_callee__ inline void asc_int82half(vector_half& dst,
85 85 
86## 调用示例86## 调用示例
87 87 
88+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
88 89 
89<!-- npu="950" id8 -->90<!-- npu="950" id8 -->
90- 
91-i代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
92- 
93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:91以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
94 92 
95```bash93```bash
96bisheng example.asc -o main --npu-arch=dav-3510 && ./main94bisheng example.asc -o main --npu-arch=dav-3510 && ./main
97```95```
96+<!-- end id8 -->
98 97 
99```c98```c
100#include <cstdint>99#include <cstdint>
@@ -190,5 +189,3 @@ int main()
190 return passed ? 0 : 1;189 return passed ? 0 : 1;
191}190}
192```191```
193- 
194-<!-- end id8 -->
@@ -85,16 +85,15 @@ __simd_callee__ inline void asc_int82int16(vector_int16_t& dst,
85 85 
86## 调用示例86## 调用示例
87 87 
88+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
88 89 
89<!-- npu="950" id8 -->90<!-- npu="950" id8 -->
90- 
91-i代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
92- 
93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:91以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
94 92 
95```bash93```bash
96bisheng example.asc -o main --npu-arch=dav-3510 && ./main94bisheng example.asc -o main --npu-arch=dav-3510 && ./main
97```95```
96+<!-- end id8 -->
98 97 
99```c98```c
100#include <cstdint>99#include <cstdint>
@@ -190,5 +189,3 @@ int main()
190 return passed ? 0 : 1;189 return passed ? 0 : 1;
191}190}
192```191```
193- 
194-<!-- end id8 -->
@@ -95,16 +95,15 @@ __simd_callee__ inline void asc_uint162uint32(vector_uint32_t& dst,
95 95 
96## 调用示例96## 调用示例
97 97 
98+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
98 99 
99<!-- npu="950" id8 -->100<!-- npu="950" id8 -->
100- 
101-u代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
102- 
103以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:101以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
104 102 
105```bash103```bash
106bisheng example.asc -o main --npu-arch=dav-3510 && ./main104bisheng example.asc -o main --npu-arch=dav-3510 && ./main
107```105```
106+<!-- end id8 -->
108 107 
109```c108```c
110#include <cstdint>109#include <cstdint>
@@ -202,5 +201,3 @@ int main()
202 return passed ? 0 : 1;201 return passed ? 0 : 1;
203}202}
204```203```
205- 
206-<!-- end id8 -->
@@ -103,16 +103,15 @@ __simd_callee__ inline void asc_uint162uint8(vector_uint8_t& dst,
103 103 
104## 调用示例104## 调用示例
105 105 
106+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
106 107 
107<!-- npu="950" id8 -->108<!-- npu="950" id8 -->
108- 
109-u代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
110- 
111以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:109以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
112 110 
113```bash111```bash
114bisheng example.asc -o main --npu-arch=dav-3510 && ./main112bisheng example.asc -o main --npu-arch=dav-3510 && ./main
115```113```
114+<!-- end id8 -->
116 115 
117```c116```c
118#include <cstdint>117#include <cstdint>
@@ -216,5 +215,3 @@ int main()
216 return passed ? 0 : 1;215 return passed ? 0 : 1;
217}216}
218```217```
219- 
220-<!-- end id8 -->
@@ -103,16 +103,15 @@ __simd_callee__ inline void asc_uint322int16(vector_int16_t& dst,
103 103 
104## 调用示例104## 调用示例
105 105 
106+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
106 107 
107<!-- npu="950" id8 -->108<!-- npu="950" id8 -->
108- 
109-u代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
110- 
111以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:109以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
112 110 
113```bash111```bash
114bisheng example.asc -o main --npu-arch=dav-3510 && ./main112bisheng example.asc -o main --npu-arch=dav-3510 && ./main
115```113```
114+<!-- end id8 -->
116 115 
117```c116```c
118#include <cstdint>117#include <cstdint>
@@ -216,5 +215,3 @@ int main()
216 return passed ? 0 : 1;215 return passed ? 0 : 1;
217}216}
218```217```
219- 
220-<!-- end id8 -->
@@ -103,16 +103,15 @@ __simd_callee__ inline void asc_uint322uint16(vector_uint16_t& dst,
103 103 
104## 调用示例104## 调用示例
105 105 
106+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
106 107 
107<!-- npu="950" id8 -->108<!-- npu="950" id8 -->
108- 
109-u代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
110- 
111以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:109以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
112 110 
113```bash111```bash
114bisheng example.asc -o main --npu-arch=dav-3510 && ./main112bisheng example.asc -o main --npu-arch=dav-3510 && ./main
115```113```
114+<!-- end id8 -->
116 115 
117```c116```c
118#include <cstdint>117#include <cstdint>
@@ -215,5 +214,3 @@ int main()
215 return passed ? 0 : 1;214 return passed ? 0 : 1;
216}215}
217```216```
218- 
219-<!-- end id8 -->
@@ -85,16 +85,15 @@ __simd_callee__ inline void asc_uint82half(vector_half& dst,
85 85 
86## 调用示例86## 调用示例
87 87 
88+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
88 89 
89<!-- npu="950" id8 -->90<!-- npu="950" id8 -->
90- 
91-u代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
92- 
93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:91以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
94 92 
95```bash93```bash
96bisheng example.asc -o main --npu-arch=dav-3510 && ./main94bisheng example.asc -o main --npu-arch=dav-3510 && ./main
97```95```
96+<!-- end id8 -->
98 97 
99```c98```c
100#include <cstdint>99#include <cstdint>
@@ -190,5 +189,3 @@ int main()
190 return passed ? 0 : 1;189 return passed ? 0 : 1;
191}190}
192```191```
193- 
194-<!-- end id8 -->
@@ -82,16 +82,15 @@ __simd_callee__ inline void asc_uint82uint16(vector_uint16_t& dst,
82 82 
83## 调用示例83## 调用示例
84 84 
85- 
86-<!-- npu="950" id8 -->
87- 
88将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。85将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
89 86 
87+<!-- npu="950" id8 -->
90以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:88以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
91 89 
92```bash90```bash
93bisheng example.asc -o main --npu-arch=dav-3510 && ./main91bisheng example.asc -o main --npu-arch=dav-3510 && ./main
94```92```
93+<!-- end id8 -->
95 94 
96 95 
97```c96```c
@@ -196,5 +195,3 @@ int main()
196 return passed ? 0 : 1;195 return passed ? 0 : 1;
197}196}
198```197```
199- 
200-<!-- end id8 -->
@@ -91,14 +91,15 @@ __simd_callee__ inline void asc_abs_sub(vector_float& dst,
91 91 
92## 调用示例92## 调用示例
93 93 
94-<!-- npu="950" id8 -->
95将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。94将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
96 95 
96+<!-- npu="950" id8 -->
97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
98 98 
99```bash99```bash
100bisheng example.asc -o main --npu-arch=dav-3510 && ./main100bisheng example.asc -o main --npu-arch=dav-3510 && ./main
101```101```
102+<!-- end id8 -->
102 103 
103```cpp104```cpp
104#include <cmath>105#include <cmath>
@@ -197,5 +198,3 @@ int main()
197 return passed ? 0 : 1;198 return passed ? 0 : 1;
198}199}
199```200```
200- 
201-<!-- end id8 -->
@@ -80,14 +80,15 @@ __simd_callee__ inline void asc_axpy(vector_float& dst,
80 80 
81## 调用示例81## 调用示例
82 82 
83-<!-- npu="950" id8 -->
84将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。83将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
85 84 
85+<!-- npu="950" id8 -->
86以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:86以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
87 87 
88```bash88```bash
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 92 
92```cpp93```cpp
93#include <cmath>94#include <cmath>
@@ -180,5 +181,3 @@ int main()
180 return passed ? 0 : 1;181 return passed ? 0 : 1;
181}182}
182```183```
183- 
184-<!-- end id8 -->
@@ -96,17 +96,18 @@
96 96 
97- 只有当输出数据类型位宽大于输入时,计算时才会有精度提升。97- 只有当输出数据类型位宽大于输入时,计算时才会有精度提升。
98 98 
99-<!-- npu="950" id8 -->
100 99 
101## 调用示例100## 调用示例
102 101 
103将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。102将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
104 103 
104+<!-- npu="950" id8 -->
105以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:105以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
106 106 
107```bash107```bash
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 111 
111```cpp112```cpp
112#include <cmath>113#include <cmath>
@@ -270,6 +271,3 @@ int main()
270 return passed ? 0 : 1;271 return passed ? 0 : 1;
271}272}
272```273```
273- 
274- 
275-<!-- end id8 -->
@@ -78,16 +78,17 @@ __simd_callee__ inline vector_float asc_fma(vector_float src0,
78- `src0``src1``src2`的数据类型需要保持一致。78- `src0``src1``src2`的数据类型需要保持一致。
79- `mask`掩码位为0时,返回值对应元素置0。79- `mask`掩码位为0时,返回值对应元素置0。
80 80 
81-<!-- npu="950" id8 -->
82## 调用示例81## 调用示例
83 82 
84将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。83将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
85 84 
85+<!-- npu="950" id8 -->
86以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:86以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
87 87 
88```bash88```bash
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 92 
92```cpp93```cpp
93#include <cmath>94#include <cmath>
@@ -196,5 +197,3 @@ int main()
196 return passed ? 0 : 1;197 return passed ? 0 : 1;
197}198}
198```199```
199- 
200-<!-- end id8 -->
@@ -81,14 +81,15 @@ __simd_callee__ inline void asc_madd(vector_float& dst,
81 81 
82## 调用示例82## 调用示例
83 83 
84-<!-- npu="950" id8 -->
85将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。84将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
86 85 
86+<!-- npu="950" id8 -->
87以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:87以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
88 88 
89```bash89```bash
90bisheng example.asc -o main --npu-arch=dav-3510 && ./main90bisheng example.asc -o main --npu-arch=dav-3510 && ./main
91```91```
92+<!-- end id8 -->
92 93 
93```cpp94```cpp
94#include <cmath>95#include <cmath>
@@ -192,5 +193,3 @@ int main()
192 return passed ? 0 : 1;193 return passed ? 0 : 1;
193}194}
194```195```
195- 
196-<!-- end id8 -->
@@ -96,16 +96,17 @@ __simd_callee__ inline void asc_mul_scalar_float2half_rn(
96- `mask`掩码位为0时,结果对应元素置0。96- `mask`掩码位为0时,结果对应元素置0。
97-`dst_pos``ASC_POSITION_EVEN`时,结果写入`dst`的偶数索引位置,奇数索引位置置0;当`dst_pos``ASC_POSITION_ODD`时,结果写入`dst`的奇数索引位置,偶数索引位置置0。97-`dst_pos``ASC_POSITION_EVEN`时,结果写入`dst`的偶数索引位置,奇数索引位置置0;当`dst_pos``ASC_POSITION_ODD`时,结果写入`dst`的奇数索引位置,偶数索引位置置0。
98 98 
99-<!-- npu="950" id8 -->
100## 调用示例99## 调用示例
101 100 
102将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。101将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
103 102 
103+<!-- npu="950" id8 -->
104以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:104以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
105 105 
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+<!-- end id8 -->
109 110 
110```cpp111```cpp
111#include <cmath>112#include <cmath>
@@ -193,5 +194,3 @@ int main()
193 return passed ? 0 : 1;194 return passed ? 0 : 1;
194}195}
195```196```
196- 
197-<!-- end id8 -->
@@ -77,16 +77,17 @@ __simd_callee__ inline void asc_mula(vector_float& dst,
77- `dst``src0``src1`的数据类型需要保持一致。77- `dst``src0``src1`的数据类型需要保持一致。
78- `mask`掩码位为0时,`dst`对应元素置0。78- `mask`掩码位为0时,`dst`对应元素置0。
79 79 
80-<!-- npu="950" id8 -->
81## 调用示例80## 调用示例
82 81 
83将以下代码保存为`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)。
84 83 
84+<!-- npu="950" id8 -->
85以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:85以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
86 86 
87```bash87```bash
88bisheng example.asc -o main --npu-arch=dav-3510 && ./main88bisheng example.asc -o main --npu-arch=dav-3510 && ./main
89```89```
90+<!-- end id8 -->
90 91 
91```cpp92```cpp
92#include <cmath>93#include <cmath>
@@ -190,5 +191,3 @@ int main()
190 return passed ? 0 : 1;191 return passed ? 0 : 1;
191}192}
192```193```
193- 
194-<!-- end id8 -->
@@ -87,14 +87,15 @@ __simd_callee__ inline void asc_relu(vector_half& dst,
87 87 
88## 调用示例88## 调用示例
89 89 
90-<!-- npu="950" id8 -->
91将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。90将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
92 91 
92+<!-- npu="950" id8 -->
93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:93以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
94 94 
95```bash95```bash
96bisheng example.asc -o main --npu-arch=dav-3510 && ./main96bisheng example.asc -o main --npu-arch=dav-3510 && ./main
97```97```
98+<!-- end id8 -->
98 99 
99```cpp100```cpp
100#include <algorithm>101#include <algorithm>
@@ -184,5 +185,3 @@ int main()
184 return passed ? 0 : 1;185 return passed ? 0 : 1;
185}186}
186```187```
187- 
188-<!-- end id8 -->
@@ -35,7 +35,7 @@ asc_xor支持两种接口:
35- 对掩码寄存器操作:35- 对掩码寄存器操作:
36 根据`mask`对两个源掩码寄存器`src0`和`src1`进行按位异或(^)操作,将结果写入目的掩码数据寄存器`dst`或通过函数返回值返回。未被`mask`筛选的位置被置为0。36 根据`mask`对两个源掩码寄存器`src0`和`src1`进行按位异或(^)操作,将结果写入目的掩码数据寄存器`dst`或通过函数返回值返回。未被`mask`筛选的位置被置为0。
37 37 
38-计算公式如下:38+伪代码如下:
39 39 
40```python40```python
41def asc_xor(dst, src0, src1, mask):41def asc_xor(dst, src0, src1, mask):
@@ -28,17 +28,49 @@
28 28 
29头文件路径为:`"c_api/reg_compute/compute/reg_permute_sel.h"`29头文件路径为:`"c_api/reg_compute/compute/reg_permute_sel.h"`
30 30 
31-将源操作数中的元素选取低8位(b16)、低16位(b32)、低32位(b64)写入目的操作数的高半部分。31+将源操作数中的元素选取低8位(b16)、低16位(b32)写入目的操作数的高半部分。
32当源操作数为vector_bool时,其按bit位进行压缩,每2bit将其中的低位bit存储至目的操作数的高半部分。32当源操作数为vector_bool时,其按bit位进行压缩,每2bit将其中的低位bit存储至目的操作数的高半部分。
33 33 
34![](../../../../figures/pack_high_draft.png)34![](../../../../figures/pack_high_draft.png)
35 35 
36## 函数原型36## 函数原型
37 37 
38-```cpp38+```c
39-__simd_callee__ inline void asc_pack_to_high(vector_uint8_t& dst, vector_uint16_t src)39+// 通过引用参数输出结果(占位符形式)
40-__simd_callee__ inline void asc_pack_to_high(vector_uint16_t& dst, vector_uint32_t src)40+__simd_callee__ inline void asc_pack_to_high(vector_<dtype_dst>& dst,
41-__simd_callee__ inline void asc_pack_to_high(vector_bool& dst, vector_bool src)41+ vector_<dtype_src> src)
42+ 
43+// 通过函数返回值返回结果(占位符形式)
44+__simd_callee__ inline vector_<dtype_dst> asc_pack_to_high(vector_<dtype_src> src)
45+ 
46+// 操作数为掩码寄存器
47+__simd_callee__ inline void asc_pack_to_high(vector_bool& dst,
48+ vector_bool src)
49+ 
50+// 操作数为掩码寄存器,通过函数返回值返回结果
51+__simd_callee__ inline vector_bool asc_pack_to_high(vector_bool src)
52+```
53+ 
54+### dtype支持数据类型
55+ 
56+dtype_src与dtype_dst支持的数据类型对如下:
57+ 
58+| dtype_src | dtype_dst |
59+|---|---|
60+| uint16_t | uint8_t |
61+| int16_t | uint8_t |
62+| uint32_t | uint16_t |
63+| int32_t | uint16_t |
64+ 
65+### 函数原型典型示例
66+ 
67+```c
68+// 示例:将uint16_t矢量数据寄存器压缩为uint8_t。
69+__simd_callee__ inline void asc_pack_to_high(vector_uint8_t& dst,
70+ vector_uint16_t src)
71+ 
72+// 示例:将uint16_t矢量数据寄存器压缩为uint8_t,通过函数返回值返回结果。
73+__simd_callee__ inline vector_uint8_t asc_pack_to_high(vector_uint16_t src)
42```74```
43 75 
44## 参数说明76## 参数说明
@@ -54,16 +86,138 @@ __simd_callee__ inline void asc_pack_to_high(vector_bool& dst, vector_bool src)
54 86 
55## 返回值说明87## 返回值说明
56 88 
57-89+- 通过引用参数输出结果的函数原型返回值,压缩结果写入`dst`。
90+- 通过函数返回值返回结果的函数原型返回压缩结果,类型为矢量数据寄存器或掩码寄存器,与`dst`的数据类型一致。
58 91 
59## 约束说明92## 约束说明
60 93 
61-94+- 本接口非AIV调用直接返回。
95+- 本接口在Vector Function(`__simd_vf__`标记的函数)内调用,源操作数和目的操作数为矢量数据寄存器或掩码寄存器。
62 96 
63## 调用示例97## 调用示例
64 98 
65-```cpp99+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
66-vector_uint8_t dst;100+ 
67-vector_uint16_t src;101+<!-- npu="950" id8 -->
68-asc_pack_to_high(dst, src); // 将src的低8位写入dst的高半部分102+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
103+ 
104+```bash
105+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
106+```
107+<!-- end id8 -->
108+ 
109+```cpp
110+#include <cstdint>
111+#include <iostream>
112+#include <vector>
113+ 
114+#include "c_api/asc_simd.h"
115+#include "acl/acl.h"
116+ 
117+namespace {
118+template <typename T>
119+void PrintData(const char* label, const std::vector<T>& values)
120+{
121+ std::cout << label << ":";
122+ const size_t count = values.size() < 8 ? values.size() : 8;
123+ for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i];
124+ if (values.size() > count) std::cout << " ...";
125+ std::cout << std::endl;
126+}
127+ 
128+template <typename T>
129+bool CompareRangeData(const std::vector<T>& actual, const std::vector<T>& expected,
130+ size_t begin, size_t count, double tolerance = 0.0)
131+{
132+ if (begin + count > actual.size() || begin + count > expected.size()) return false;
133+ for (size_t i = begin; i < begin + count; ++i) {
134+ if (actual[i] == expected[i]) continue;
135+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
136+ if (diff > tolerance || diff < -tolerance) return false;
137+ }
138+ return true;
139+}
140+ 
141+constexpr uint32_t SRC_COUNT = 128;
142+constexpr uint32_t DST_COUNT = 256;
143+ 
144+__simd_vf__ inline void Pack(__ubuf__ uint8_t* low, __ubuf__ uint8_t* high, __ubuf__ uint16_t* src)
145+{
146+ vector_uint8_t low_reg;
147+ vector_uint8_t high_reg;
148+ vector_uint16_t src_reg;
149+ uint32_t count = DST_COUNT;
150+ vector_bool dst_mask = asc_update_mask_b8(count);
151+ asc_loadalign(src_reg, src);
152+ asc_duplicate_scalar(low_reg, static_cast<uint8_t>(0), dst_mask);
153+ asc_duplicate_scalar(high_reg, static_cast<uint8_t>(0), dst_mask);
154+
155+ asc_pack_to_high(high_reg, src_reg);
156+
157+ asc_storealign(high, high_reg, dst_mask);
158+}
159+ 
160+__global__ __vector__ void AscPackToHighKernel(__gm__ uint8_t* low, __gm__ uint8_t* high, __gm__ uint16_t* src)
161+{
162+ asc_init();
163+ __ubuf__ uint8_t low_local[DST_COUNT];
164+ __ubuf__ uint8_t high_local[DST_COUNT];
165+ __ubuf__ uint16_t src_local[SRC_COUNT];
166+ asc_copy_gm2ub_align(src_local, src, SRC_COUNT * sizeof(uint16_t));
167+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
168+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
169+ Pack(low_local, high_local, src_local);
170+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
171+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
172+ asc_copy_ub2gm_align(low, low_local, DST_COUNT);
173+ asc_copy_ub2gm_align(high, high_local, DST_COUNT);
174+ asc_sync();
175+}
176+} // namespace
177+ 
178+int main()
179+{
180+ std::vector<uint16_t> src(SRC_COUNT);
181+ std::vector<uint8_t> low(DST_COUNT, 0);
182+ std::vector<uint8_t> high(DST_COUNT, 0);
183+ std::vector<uint8_t> golden_low(DST_COUNT, 0);
184+ std::vector<uint8_t> golden_high(DST_COUNT, 0);
185+ for (uint32_t i = 0; i < SRC_COUNT; ++i) {
186+ src[i] = static_cast<uint16_t>(0x1200U + i);
187+
188+ golden_high[SRC_COUNT + i] = static_cast<uint8_t>(src[i]);
189+ }
190+ aclInit(nullptr);
191+ aclrtSetDevice(0);
192+ uint16_t* src_device = nullptr;
193+ aclrtMalloc(reinterpret_cast<void**>(&src_device), (SRC_COUNT) * sizeof(uint16_t),
194+ ACL_MEM_MALLOC_HUGE_FIRST);
195+ uint8_t* low_device = nullptr;
196+ aclrtMalloc(reinterpret_cast<void**>(&low_device), (DST_COUNT) * sizeof(uint8_t),
197+ ACL_MEM_MALLOC_HUGE_FIRST);
198+ uint8_t* high_device = nullptr;
199+ aclrtMalloc(reinterpret_cast<void**>(&high_device), (DST_COUNT) * sizeof(uint8_t),
200+ ACL_MEM_MALLOC_HUGE_FIRST);
201+ aclrtMemcpy(src_device, src.size() * sizeof(uint16_t), src.data(), src.size() * sizeof(uint16_t),
202+ ACL_MEMCPY_HOST_TO_DEVICE);
203+ AscPackToHighKernel<<<1, 0>>>(low_device, high_device, src_device);
204+ aclrtSynchronizeDevice();
205+ aclrtMemcpy(low.data(), low.size() * sizeof(uint8_t), low_device, low.size() * sizeof(uint8_t),
206+ ACL_MEMCPY_DEVICE_TO_HOST);
207+ aclrtMemcpy(high.data(), high.size() * sizeof(uint8_t), high_device, high.size() * sizeof(uint8_t),
208+ ACL_MEMCPY_DEVICE_TO_HOST);
209+ PrintData("Input", src);
210+ PrintData("Low output", low);
211+ PrintData("Low Golden", golden_low);
212+ PrintData("High output", high);
213+ PrintData("High Golden", golden_high);
214+ const bool passed = CompareRangeData(high, golden_high, SRC_COUNT, SRC_COUNT);
215+ std::cout << (passed ? "[Success] asc_pack_to_high passed." : "[Failed] asc_pack_to_high failed.") << std::endl;
216+ aclrtFree(high_device);
217+ aclrtFree(low_device);
218+ aclrtFree(src_device);
219+ aclrtResetDevice(0);
220+ aclFinalize();
221+ return passed ? 0 : 1;
222+}
69```223```
@@ -28,17 +28,49 @@
28 28 
29头文件路径为:`"c_api/reg_compute/compute/reg_permute_sel.h"`29头文件路径为:`"c_api/reg_compute/compute/reg_permute_sel.h"`
30 30 
31-将源操作数中的元素选取偶数位bit(bool)、低8位(b16)、低16位(b32)、低32位(b64)写入目的操作数的低半部分。31+将源操作数中的元素选取偶数位bit(bool)、低8位(b16)、低16位(b32)写入目的操作数的低半部分。
32当源操作数为vector_bool时,其按bit位进行压缩,每2bit将其中的低位bit存储至目的操作数的低半部分。32当源操作数为vector_bool时,其按bit位进行压缩,每2bit将其中的低位bit存储至目的操作数的低半部分。
33 33 
34![](../../../../figures/pack_low_draft.png)34![](../../../../figures/pack_low_draft.png)
35 35 
36## 函数原型36## 函数原型
37 37 
38-```cpp38+```c
39-__simd_callee__ inline void asc_pack_to_low(vector_uint8_t& dst, vector_uint16_t src)39+// 通过引用参数输出结果(占位符形式)
40-__simd_callee__ inline void asc_pack_to_low(vector_uint16_t& dst, vector_uint32_t src)40+__simd_callee__ inline void asc_pack_to_low(vector_<dtype_dst>& dst,
41-__simd_callee__ inline void asc_pack_to_low(vector_bool& dst, vector_bool src)41+ vector_<dtype_src> src)
42+ 
43+// 通过函数返回值返回结果(占位符形式)
44+__simd_callee__ inline vector_<dtype_dst> asc_pack_to_low(vector_<dtype_src> src)
45+ 
46+// 操作数为掩码寄存器
47+__simd_callee__ inline void asc_pack_to_low(vector_bool& dst,
48+ vector_bool src)
49+ 
50+// 操作数为掩码寄存器,通过函数返回值返回结果
51+__simd_callee__ inline vector_bool asc_pack_to_low(vector_bool src)
52+```
53+ 
54+### dtype支持数据类型
55+ 
56+dtype_src与dtype_dst支持的数据类型对如下:
57+ 
58+| dtype_src | dtype_dst |
59+|---|---|
60+| uint16_t | uint8_t |
61+| int16_t | uint8_t |
62+| uint32_t | uint16_t |
63+| int32_t | uint16_t |
64+ 
65+### 函数原型典型示例
66+ 
67+```c
68+// 示例:将uint16_t矢量数据寄存器压缩为uint8_t。
69+__simd_callee__ inline void asc_pack_to_low(vector_uint8_t& dst,
70+ vector_uint16_t src)
71+ 
72+// 示例:将uint16_t矢量数据寄存器压缩为uint8_t,通过函数返回值返回结果。
73+__simd_callee__ inline vector_uint8_t asc_pack_to_low(vector_uint16_t src)
42```74```
43 75 
44## 参数说明76## 参数说明
@@ -54,16 +86,138 @@ __simd_callee__ inline void asc_pack_to_low(vector_bool& dst, vector_bool src)
54 86 
55## 返回值说明87## 返回值说明
56 88 
57-89+- 通过引用参数输出结果的函数原型返回值,压缩结果写入`dst`。
90+- 通过函数返回值返回结果的函数原型返回压缩结果,类型为矢量数据寄存器或掩码寄存器,与`dst`的数据类型一致。
58 91 
59## 约束说明92## 约束说明
60 93 
61-94+- 本接口非AIV调用直接返回。
95+- 本接口在Vector Function(`__simd_vf__`标记的函数)内调用,源操作数和目的操作数为矢量数据寄存器或掩码寄存器。
62 96 
63## 调用示例97## 调用示例
64 98 
65-```cpp99+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
66-vector_uint8_t dst;100+ 
67-vector_uint16_t src;101+<!-- npu="950" id8 -->
68-asc_pack_to_low(dst, src); // 将src的低8位写入dst的低半部分102+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
103+ 
104+```bash
105+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
106+```
107+<!-- end id8 -->
108+ 
109+```cpp
110+#include <cstdint>
111+#include <iostream>
112+#include <vector>
113+ 
114+#include "c_api/asc_simd.h"
115+#include "acl/acl.h"
116+ 
117+namespace {
118+template <typename T>
119+void PrintData(const char* label, const std::vector<T>& values)
120+{
121+ std::cout << label << ":";
122+ const size_t count = values.size() < 8 ? values.size() : 8;
123+ for (size_t i = 0; i < count; ++i) std::cout << ' ' << +values[i];
124+ if (values.size() > count) std::cout << " ...";
125+ std::cout << std::endl;
126+}
127+ 
128+template <typename T>
129+bool CompareRangeData(const std::vector<T>& actual, const std::vector<T>& expected,
130+ size_t begin, size_t count, double tolerance = 0.0)
131+{
132+ if (begin + count > actual.size() || begin + count > expected.size()) return false;
133+ for (size_t i = begin; i < begin + count; ++i) {
134+ if (actual[i] == expected[i]) continue;
135+ const double diff = static_cast<double>(actual[i]) - static_cast<double>(expected[i]);
136+ if (diff > tolerance || diff < -tolerance) return false;
137+ }
138+ return true;
139+}
140+ 
141+constexpr uint32_t SRC_COUNT = 128;
142+constexpr uint32_t DST_COUNT = 256;
143+ 
144+__simd_vf__ inline void Pack(__ubuf__ uint8_t* low, __ubuf__ uint8_t* high, __ubuf__ uint16_t* src)
145+{
146+ vector_uint8_t low_reg;
147+ vector_uint8_t high_reg;
148+ vector_uint16_t src_reg;
149+ uint32_t count = DST_COUNT;
150+ vector_bool dst_mask = asc_update_mask_b8(count);
151+ asc_loadalign(src_reg, src);
152+ asc_duplicate_scalar(low_reg, static_cast<uint8_t>(0), dst_mask);
153+ asc_duplicate_scalar(high_reg, static_cast<uint8_t>(0), dst_mask);
154+ asc_pack_to_low(low_reg, src_reg);
155+
156+ asc_storealign(low, low_reg, dst_mask);
157+
158+}
159+ 
160+__global__ __vector__ void AscPackToLowKernel(__gm__ uint8_t* low, __gm__ uint8_t* high, __gm__ uint16_t* src)
161+{
162+ asc_init();
163+ __ubuf__ uint8_t low_local[DST_COUNT];
164+ __ubuf__ uint8_t high_local[DST_COUNT];
165+ __ubuf__ uint16_t src_local[SRC_COUNT];
166+ asc_copy_gm2ub_align(src_local, src, SRC_COUNT * sizeof(uint16_t));
167+ asc_sync_notify(PIPE_MTE2, PIPE_V, EVENT_ID0);
168+ asc_sync_wait(PIPE_MTE2, PIPE_V, EVENT_ID0);
169+ Pack(low_local, high_local, src_local);
170+ asc_sync_notify(PIPE_V, PIPE_MTE3, EVENT_ID0);
171+ asc_sync_wait(PIPE_V, PIPE_MTE3, EVENT_ID0);
172+ asc_copy_ub2gm_align(low, low_local, DST_COUNT);
173+ asc_copy_ub2gm_align(high, high_local, DST_COUNT);
174+ asc_sync();
175+}
176+} // namespace
177+ 
178+int main()
179+{
180+ std::vector<uint16_t> src(SRC_COUNT);
181+ std::vector<uint8_t> low(DST_COUNT, 0);
182+ std::vector<uint8_t> high(DST_COUNT, 0);
183+ std::vector<uint8_t> golden_low(DST_COUNT, 0);
184+ std::vector<uint8_t> golden_high(DST_COUNT, 0);
185+ for (uint32_t i = 0; i < SRC_COUNT; ++i) {
186+ src[i] = static_cast<uint16_t>(0x1200U + i);
187+ golden_low[i] = static_cast<uint8_t>(src[i]);
188+
189+ }
190+ aclInit(nullptr);
191+ aclrtSetDevice(0);
192+ uint16_t* src_device = nullptr;
193+ aclrtMalloc(reinterpret_cast<void**>(&src_device), (SRC_COUNT) * sizeof(uint16_t),
194+ ACL_MEM_MALLOC_HUGE_FIRST);
195+ uint8_t* low_device = nullptr;
196+ aclrtMalloc(reinterpret_cast<void**>(&low_device), (DST_COUNT) * sizeof(uint8_t),
197+ ACL_MEM_MALLOC_HUGE_FIRST);
198+ uint8_t* high_device = nullptr;
199+ aclrtMalloc(reinterpret_cast<void**>(&high_device), (DST_COUNT) * sizeof(uint8_t),
200+ ACL_MEM_MALLOC_HUGE_FIRST);
201+ aclrtMemcpy(src_device, src.size() * sizeof(uint16_t), src.data(), src.size() * sizeof(uint16_t),
202+ ACL_MEMCPY_HOST_TO_DEVICE);
203+ AscPackToLowKernel<<<1, 0>>>(low_device, high_device, src_device);
204+ aclrtSynchronizeDevice();
205+ aclrtMemcpy(low.data(), low.size() * sizeof(uint8_t), low_device, low.size() * sizeof(uint8_t),
206+ ACL_MEMCPY_DEVICE_TO_HOST);
207+ aclrtMemcpy(high.data(), high.size() * sizeof(uint8_t), high_device, high.size() * sizeof(uint8_t),
208+ ACL_MEMCPY_DEVICE_TO_HOST);
209+ PrintData("Input", src);
210+ PrintData("Low output", low);
211+ PrintData("Low Golden", golden_low);
212+ PrintData("High output", high);
213+ PrintData("High Golden", golden_high);
214+ const bool passed = CompareRangeData(low, golden_low, 0, SRC_COUNT);
215+ std::cout << (passed ? "[Success] asc_pack_to_low passed." : "[Failed] asc_pack_to_low failed.") << std::endl;
216+ aclrtFree(high_device);
217+ aclrtFree(low_device);
218+ aclrtFree(src_device);
219+ aclrtResetDevice(0);
220+ aclFinalize();
221+ return passed ? 0 : 1;
222+}
69```223```
@@ -93,15 +93,15 @@ __simd_callee__ inline void asc_select(vector_half& dst,
93 93 
94## 调用示例94## 调用示例
95 95 
96-<!-- npu="950" id8 -->
97- 
98将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。96将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
99 97 
98+<!-- npu="950" id8 -->
100以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:99以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
101 100 
102```bash101```bash
103bisheng example.asc -o main --npu-arch=dav-3510 && ./main102bisheng example.asc -o main --npu-arch=dav-3510 && ./main
104```103```
104+<!-- end id8 -->
105 105 
106```c106```c
107#include <cstdint>107#include <cstdint>
@@ -200,4 +200,3 @@ int main()
200}200}
201```201```
202 202 
203-<!-- end id8 -->
@@ -97,14 +97,15 @@ __simd_callee__ inline void asc_pair_reduce_sum(vector_half& dst,
97 97 
98## 调用示例98## 调用示例
99 99 
100-<!-- npu="950" id8 -->
101将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。100将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
102 101 
102+<!-- npu="950" id8 -->
103以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:103以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
104 104 
105```bash105```bash
106bisheng example.asc -o main --npu-arch=dav-3510 && ./main106bisheng example.asc -o main --npu-arch=dav-3510 && ./main
107```107```
108+<!-- end id8 -->
108 109 
109```cpp110```cpp
110#include <cmath>111#include <cmath>
@@ -191,5 +192,3 @@ int main()
191 return passed ? 0 : 1;192 return passed ? 0 : 1;
192}193}
193```194```
194- 
195-<!-- end id8 -->
@@ -97,14 +97,15 @@ __simd_callee__ inline void asc_reduce_max(vector_int16_t& dst,
97 97 
98## 调用示例98## 调用示例
99 99 
100-<!-- npu="950" id8 -->
101将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。100将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
102 101 
102+<!-- npu="950" id8 -->
103以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:103以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
104 104 
105```bash105```bash
106bisheng example.asc -o main --npu-arch=dav-3510 && ./main106bisheng example.asc -o main --npu-arch=dav-3510 && ./main
107```107```
108+<!-- end id8 -->
108 109 
109```cpp110```cpp
110// 本例使用half类型演示归约结果"值+索引"的分别提取,111// 本例使用half类型演示归约结果"值+索引"的分别提取,
@@ -170,5 +171,3 @@ int main()
170 return passed ? 0 : 1;171 return passed ? 0 : 1;
171}172}
172```173```
173- 
174-<!-- end id8 -->
@@ -92,14 +92,15 @@ __simd_callee__ inline void asc_reduce_max_datablock(vector_half& dst,
92 92 
93## 调用示例93## 调用示例
94 94 
95-<!-- npu="950" id8 -->
96将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。95将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
97 96 
97+<!-- npu="950" id8 -->
98以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:98以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
99 99 
100```bash100```bash
101bisheng example.asc -o main --npu-arch=dav-3510 && ./main101bisheng example.asc -o main --npu-arch=dav-3510 && ./main
102```102```
103+<!-- end id8 -->
103 104 
104```cpp105```cpp
105#include <cstdint>106#include <cstdint>
@@ -172,5 +173,3 @@ int main()
172 return passed ? 0 : 1;173 return passed ? 0 : 1;
173}174}
174```175```
175- 
176-<!-- end id8 -->
@@ -97,14 +97,15 @@ __simd_callee__ inline void asc_reduce_min(vector_half& dst,
97 97 
98## 调用示例98## 调用示例
99 99 
100-<!-- npu="950" id8 -->
101将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。100将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
102 101 
102+<!-- npu="950" id8 -->
103以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:103以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
104 104 
105```bash105```bash
106bisheng example.asc -o main --npu-arch=dav-3510 && ./main106bisheng example.asc -o main --npu-arch=dav-3510 && ./main
107```107```
108+<!-- end id8 -->
108 109 
109```cpp110```cpp
110// 本例使用half类型演示归约结果"值+索引"的分别提取,111// 本例使用half类型演示归约结果"值+索引"的分别提取,
@@ -170,5 +171,3 @@ int main()
170 return passed ? 0 : 1;171 return passed ? 0 : 1;
171}172}
172```173```
173- 
174-<!-- end id8 -->
@@ -91,14 +91,15 @@ __simd_callee__ inline void asc_reduce_min_datablock(vector_half& dst,
91 91 
92## 调用示例92## 调用示例
93 93 
94-<!-- npu="950" id8 -->
95将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。94将以下代码保存为`example.asc`后,可通过`bisheng`命令编译运行。其中,`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
96 95 
96+<!-- npu="950" id8 -->
97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:97以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
98 98 
99```bash99```bash
100bisheng example.asc -o main --npu-arch=dav-3510 && ./main100bisheng example.asc -o main --npu-arch=dav-3510 && ./main
101```101```
102+<!-- end id8 -->
102 103 
103```cpp104```cpp
104#include <cstdint>105#include <cstdint>
@@ -171,5 +172,3 @@ int main()
171 return passed ? 0 : 1;172 return passed ? 0 : 1;
172}173}
173```174```
174- 
175-<!-- end id8 -->
@@ -118,14 +118,15 @@ __simd_callee__ inline void asc_reduce_sum(vector_float& dst,
118 118 
119## 调用示例119## 调用示例
120 120 
121-<!-- npu="950" id8 -->
122将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。121将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
123 122 
123+<!-- npu="950" id8 -->
124以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:124以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
125 125 
126```bash126```bash
127bisheng example.asc -o main --npu-arch=dav-3510 && ./main127bisheng example.asc -o main --npu-arch=dav-3510 && ./main
128```128```
129+<!-- end id8 -->
129 130 
130```c131```c
131#include <cstdint>132#include <cstdint>
@@ -200,5 +201,3 @@ int main()
200 return passed ? 0 : 1;201 return passed ? 0 : 1;
201}202}
202```203```
203- 
204-<!-- end id8 -->
@@ -124,15 +124,15 @@ __simd_callee__ inline void asc_reduce_sum_datablock(vector_half& dst,
124 124 
125## 调用示例125## 调用示例
126 126 
127-<!-- npu="950" id8 -->
128- 
129将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。127将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[__NPU_ARCH__](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
130 128 
129+<!-- npu="950" id8 -->
131以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:130以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
132 131 
133```bash132```bash
134bisheng example.asc -o main --npu-arch=dav-3510 && ./main133bisheng example.asc -o main --npu-arch=dav-3510 && ./main
135```134```
135+<!-- end id8 -->
136 136 
137```c137```c
138#include <cstdint>138#include <cstdint>
@@ -208,4 +208,3 @@ int main()
208}208}
209```209```
210 210 
211-<!-- end id8 -->
@@ -65,7 +65,6 @@ PIPE_S
65将以下代码保存为`example.asc`后,执行对应的编译运行命令。65将以下代码保存为`example.asc`后,执行对应的编译运行命令。
66 66 
67<!-- npu="950" id8 -->67<!-- npu="950" id8 -->
68- 
69以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:68以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
70 69 
71```bash70```bash
@@ -32,7 +32,7 @@
32 32 
33## 函数原型33## 函数原型
34 34 
35-```cpp35+```c
36__aicore__ inline int64_t asc_get_ctrl()36__aicore__ inline int64_t asc_get_ctrl()
37```37```
38 38 
@@ -54,4 +54,77 @@ PIPE_S
54 54 
55## 调用示例55## 调用示例
56 56 
57-57+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
58+ 
59+<!-- npu="950" id8 -->
60+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
61+ 
62+```bash
63+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
64+```
65+<!-- end id8 -->
66+ 
67+```cpp
68+#include <cstdint>
69+#include <iomanip>
70+#include <iostream>
71+#include <vector>
72+#include "c_api/asc_simd.h"
73+#include "acl/acl.h"
74+ 
75+namespace {
76+constexpr uint32_t ELEMENTS = 16;
77+constexpr uint32_t BYTES = ELEMENTS * sizeof(uint64_t);
78+ 
79+void print_data(const char* label, const std::vector<uint64_t>& data)
80+{
81+ std::cout << label << ":";
82+ for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << data[i];
83+ std::cout << " ..." << std::endl;
84+}
85+ 
86+void print_hex(const char* label, const std::vector<uint64_t>& data)
87+{
88+ std::cout << label << ":" << std::hex;
89+ for (uint32_t i = 0; i < 4; ++i) std::cout << " 0x" << data[i];
90+ std::cout << std::dec << std::endl;
91+}
92+ 
93+__global__ __vector__ void asc_get_ctrl_kernel(__gm__ uint64_t* output)
94+{
95+ asc_init();
96+ const uint64_t before = static_cast<uint64_t>(asc_get_ctrl());
97+ const uint64_t changed = before ^ (1ULL << 53);
98+ asc_set_ctrl(changed);
99+ const uint64_t after = static_cast<uint64_t>(asc_get_ctrl());
100+ asc_set_ctrl(before);
101+ const uint64_t restored = static_cast<uint64_t>(asc_get_ctrl());
102+ output[0] = before;
103+ output[1] = after;
104+ output[2] = restored;
105+ asc_dcci_single(output);
106+}
107+ 
108+} // namespace
109+ 
110+int main()
111+{
112+ std::vector<uint64_t> output(ELEMENTS, ~0ULL);
113+ aclInit(nullptr);
114+ aclrtSetDevice(0);
115+ uint64_t* output_device = nullptr;
116+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
117+ aclrtMemcpy(output_device, BYTES, output.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
118+ asc_get_ctrl_kernel<<<1, 0>>>(output_device);
119+ aclrtSynchronizeDevice();
120+ aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
121+ print_data("Observed", output);
122+ print_hex("Observed hex", output);
123+ const bool passed = ((output[0] ^ output[1]) & (1ULL << 53)) != 0 && output[2] == output[0];
124+ std::cout << (passed ? "[Success] asc_get_ctrl passed." : "[Failed] asc_get_ctrl failed.") << std::endl;
125+ aclrtFree(output_device);
126+ aclrtResetDevice(0);
127+ aclFinalize();
128+ return passed ? 0 : 1;
129+}
130+```
@@ -32,7 +32,7 @@
32 32 
33## 函数原型33## 函数原型
34 34 
35-```cpp35+```c
36__aicore__ inline void asc_set_ctrl(uint64_t config)36__aicore__ inline void asc_set_ctrl(uint64_t config)
37```37```
38 38 
@@ -71,10 +71,77 @@ PIPE_S
71 71 
72## 调用示例72## 调用示例
73 73 
74-```cpp74+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
75-// 将CTRL[48]设置为非饱和模式75+ 
76-uint64_t ori_config = asc_get_ctrl();76+<!-- npu="950" id8 -->
77-uint64_t mask = (static_cast<uint64_t>(1) << 48);77+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
78-uint64_t config = ori_config | mask;78+ 
79-asc_set_ctrl(config);79+```bash
80+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
81+```
82+<!-- end id8 -->
83+ 
84+```cpp
85+#include <cstdint>
86+#include <iomanip>
87+#include <iostream>
88+#include <vector>
89+#include "c_api/asc_simd.h"
90+#include "acl/acl.h"
91+ 
92+namespace {
93+constexpr uint32_t ELEMENTS = 16;
94+constexpr uint32_t BYTES = ELEMENTS * sizeof(uint64_t);
95+ 
96+void print_data(const char* label, const std::vector<uint64_t>& data)
97+{
98+ std::cout << label << ":";
99+ for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << data[i];
100+ std::cout << " ..." << std::endl;
101+}
102+ 
103+void print_hex(const char* label, const std::vector<uint64_t>& data)
104+{
105+ std::cout << label << ":" << std::hex;
106+ for (uint32_t i = 0; i < 4; ++i) std::cout << " 0x" << data[i];
107+ std::cout << std::dec << std::endl;
108+}
109+ 
110+__global__ __vector__ void asc_set_ctrl_kernel(__gm__ uint64_t* output)
111+{
112+ asc_init();
113+ const uint64_t before = static_cast<uint64_t>(asc_get_ctrl());
114+ const uint64_t changed = before ^ (1ULL << 56);
115+ asc_set_ctrl(changed);
116+ const uint64_t after = static_cast<uint64_t>(asc_get_ctrl());
117+ asc_set_ctrl(before);
118+ const uint64_t restored = static_cast<uint64_t>(asc_get_ctrl());
119+ output[0] = before;
120+ output[1] = after;
121+ output[2] = restored;
122+ asc_dcci_single(output);
123+}
124+ 
125+} // namespace
126+ 
127+int main()
128+{
129+ std::vector<uint64_t> output(ELEMENTS, ~0ULL);
130+ aclInit(nullptr);
131+ aclrtSetDevice(0);
132+ uint64_t* output_device = nullptr;
133+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
134+ aclrtMemcpy(output_device, BYTES, output.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
135+ asc_set_ctrl_kernel<<<1, 0>>>(output_device);
136+ aclrtSynchronizeDevice();
137+ aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
138+ print_data("Observed", output);
139+ print_hex("Observed hex", output);
140+ const bool passed = ((output[0] ^ output[1]) & (1ULL << 56)) != 0 && output[2] == output[0];
141+ std::cout << (passed ? "[Success] asc_set_ctrl passed." : "[Failed] asc_set_ctrl failed.") << std::endl;
142+ aclrtFree(output_device);
143+ aclrtResetDevice(0);
144+ aclFinalize();
145+ return passed ? 0 : 1;
146+}
80```147```
@@ -5,7 +5,6 @@
5AI Core内部的执行单元(如MTE2搬运单元、Vector计算单元等)以异步并行的方式运行,在读写同一存储资源时可能存在数据依赖关系。为确保数据一致性及计算正确性,需通过同步控制协调操作时序。5AI Core内部的执行单元(如MTE2搬运单元、Vector计算单元等)以异步并行的方式运行,在读写同一存储资源时可能存在数据依赖关系。为确保数据一致性及计算正确性,需通过同步控制协调操作时序。
6 6 
7<!-- npu="950" id2 -->7<!-- npu="950" id2 -->
8- 
9针对[NPU架构3510](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch),硬件架构图如下,高亮部分展示了并行执行的计算单元和搬运单元。8针对[NPU架构3510](../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch),硬件架构图如下,高亮部分展示了并行执行的计算单元和搬运单元。
10 9 
11**图1** NPU架构3510架构图 10**图1** NPU架构3510架构图
@@ -28,11 +28,11 @@
28 28 
29头文件路径为:`"c_api/utils/sys_var.h"`29头文件路径为:`"c_api/utils/sys_var.h"`
30 30 
31-取状态信息31+当前核上下文中STATUS特殊寄存器的64bit整体值并按`int64_t`返回。STATUS是64bit状态寄存器,记录核运行过程中各类告警事件的发生情况,各状态位一旦因告警行为置位即保持,直到复位或被读取清除
32 32 
33## 函数原型33## 函数原型
34 34 
35-```cpp35+```c
36__aicore__ inline int64_t asc_get_status()36__aicore__ inline int64_t asc_get_status()
GRJ_XIDUO
GRJ_XIDUOGRJ_XIDUO14 天前

返回值说明部分,表的标题和文字渲染在一行,如下所示,预期应当换行: “状态信息。各bit含义如下: 表1 状态寄存器比特位说明”

likedislike
37```37```
38 38 
@@ -43,6 +43,7 @@ __aicore__ inline int64_t asc_get_status()
43## 返回值说明43## 返回值说明
44 44 
45状态信息。各bit含义如下:45状态信息。各bit含义如下:
46+ 
46**表1** 状态寄存器比特位说明47**表1** 状态寄存器比特位说明
47 48 
48| bit范围 | 含义 |49| bit范围 | 含义 |
@@ -66,11 +67,75 @@ PIPE_S
66 67 
67## 约束说明68## 约束说明
68 69 
69-70+- 本接口为只读查询接口,仅读取状态寄存器当前值,可在AIC与AIV中调用。
71+- 本接口不支持纯SIMT编译模式。
70 72 
71## 调用示例73## 调用示例
72 74 
73-```cpp75+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
74-int64_t status = asc_get_status();76+ 
75-printf("status is %x", status);// 需用%x将其打印成十六进制的数77+<!-- npu="950" id8 -->
78+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
79+ 
80+```bash
81+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
82+```
83+<!-- end id8 -->
84+ 
85+```cpp
86+#include <cstdint>
87+#include <iomanip>
88+#include <iostream>
89+#include <vector>
90+#include "c_api/asc_simd.h"
91+#include "acl/acl.h"
92+ 
93+namespace {
94+constexpr uint32_t ELEMENTS = 16;
95+constexpr uint32_t BYTES = ELEMENTS * sizeof(uint64_t);
96+ 
97+void print_data(const char* label, const std::vector<uint64_t>& data)
98+{
99+ std::cout << label << ":";
100+ for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << data[i];
101+ std::cout << " ..." << std::endl;
102+}
103+ 
104+void print_hex(const char* label, const std::vector<uint64_t>& data)
105+{
106+ std::cout << label << ":" << std::hex;
107+ for (uint32_t i = 0; i < 4; ++i) std::cout << " 0x" << data[i];
108+ std::cout << std::dec << std::endl;
109+}
110+ 
111+__global__ __vector__ void asc_get_status_kernel(__gm__ uint64_t* output)
112+{
113+ asc_init();
114+ output[0] = static_cast<uint64_t>(asc_get_status());
115+ output[1] = static_cast<uint64_t>(asc_get_status());
116+ asc_dcci_single(output);
117+}
118+ 
119+} // namespace
120+ 
121+int main()
122+{
123+ std::vector<uint64_t> output(ELEMENTS, ~0ULL);
124+ aclInit(nullptr);
125+ aclrtSetDevice(0);
126+ uint64_t* output_device = nullptr;
127+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
128+ aclrtMemcpy(output_device, BYTES, output.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
129+ asc_get_status_kernel<<<1, 0>>>(output_device);
130+ aclrtSynchronizeDevice();
131+ aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
132+ print_data("Observed", output);
133+ print_hex("Observed hex", output);
134+ const bool passed = output[0] == output[1];
135+ std::cout << (passed ? "[Success] asc_get_status passed." : "[Failed] asc_get_status failed.") << std::endl;
136+ aclrtFree(output_device);
137+ aclrtResetDevice(0);
138+ aclFinalize();
139+ return passed ? 0 : 1;
140+}
76```141```
@@ -28,11 +28,11 @@
28 28 
29头文件路径为:`"c_api/utils/sys_var.h"`29头文件路径为:`"c_api/utils/sys_var.h"`
30 30 
31-获取AI Core上VectorID31+在[分离模式](../../../../../guide/programming_guide/advanced_programming/hardware_implementation/basic_architecture.md)下获取逻辑AI Core上Cube Core(AIC)或者Vector Core(AIV)逻辑索引
32 32 
33## 函数原型33## 函数原型
34 34 
35-```cpp35+```c
36__aicore__ inline int64_t asc_get_sub_block_id()36__aicore__ inline int64_t asc_get_sub_block_id()
37```37```
38 38 
@@ -42,7 +42,9 @@ __aicore__ inline int64_t asc_get_sub_block_id()
42 42 
43## 返回值说明43## 返回值说明
44 44 
45-返回Vector核ID45+返回逻辑AI Core上Cube Core(AIC)或者Vector Core(AIV)的逻辑索引
46+ 
47+SubBlockIdx的取值范围为[0, SubBlockNum),SubBlockNum可以通过[GetSubBlockNum](asc_get_sub_block_num.md)获取。
46 48 
47## 流水类型49## 流水类型
48 50 
@@ -50,8 +52,77 @@ PIPE_S
50 52 
51## 约束说明53## 约束说明
52 54 
53-55+本接口为只读查询接口,不修改任何寄存器或存储状态。
54 56 
55## 调用示例57## 调用示例
56 58 
57-59+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
60+ 
61+<!-- npu="950" id8 -->
62+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
63+ 
64+```bash
65+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
66+```
67+<!-- end id8 -->
68+ 
69+```cpp
70+#include <cstdint>
71+#include <iomanip>
72+#include <iostream>
73+#include <vector>
74+#include "c_api/asc_simd.h"
75+#include "acl/acl.h"
76+ 
77+namespace {
78+constexpr uint32_t ELEMENTS = 16;
79+constexpr uint32_t BYTES = ELEMENTS * sizeof(uint64_t);
80+ 
81+void print_data(const char* label, const std::vector<uint64_t>& data)
82+{
83+ std::cout << label << ":";
84+ for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << data[i];
85+ std::cout << " ..." << std::endl;
86+}
87+ 
88+void print_hex(const char* label, const std::vector<uint64_t>& data)
89+{
90+ std::cout << label << ":" << std::hex;
91+ for (uint32_t i = 0; i < 4; ++i) std::cout << " 0x" << data[i];
92+ std::cout << std::dec << std::endl;
93+}
94+ 
95+__global__ __vector__ void asc_get_sub_block_id_kernel(__gm__ uint64_t* output)
96+{
97+ asc_init();
98+ const uint64_t sub_id = static_cast<uint64_t>(asc_get_sub_block_id());
99+ const uint64_t sub_num = static_cast<uint64_t>(asc_get_sub_block_num());
100+ output[0] = sub_id;
101+ output[1] = sub_id;
102+ output[2] = sub_num;
103+ asc_dcci_single(output);
104+}
105+ 
106+} // namespace
107+ 
108+int main()
109+{
110+ std::vector<uint64_t> output(ELEMENTS, ~0ULL);
111+ aclInit(nullptr);
112+ aclrtSetDevice(0);
113+ uint64_t* output_device = nullptr;
114+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
115+ aclrtMemcpy(output_device, BYTES, output.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
116+ asc_get_sub_block_id_kernel<<<1, 0>>>(output_device);
117+ aclrtSynchronizeDevice();
118+ aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
119+ print_data("Observed", output);
120+ print_hex("Observed hex", output);
121+ const bool passed = output[2] > 0 && output[1] < output[2];
122+ std::cout << (passed ? "[Success] asc_get_sub_block_id passed." : "[Failed] asc_get_sub_block_id failed.") << std::endl;
123+ aclrtFree(output_device);
124+ aclrtResetDevice(0);
125+ aclFinalize();
126+ return passed ? 0 : 1;
127+}
128+```
@@ -28,11 +28,11 @@
28 28 
29头文件路径为:`"c_api/utils/sys_var.h"`29头文件路径为:`"c_api/utils/sys_var.h"`
30 30 
31-分离模式下,获一个AI Core上Cube Core(AIC)或Vector Core(AIV)的数量。31+在[分离模式](../../../../../guide/programming_guide/advanced_programming/hardware_implementation/basic_architecture.md)当前AI Core上Cube Core(AIC)或Vector Core(AIV)的数量。该值在kernel启动前配置,运行中不可修改。
32 32 
33## 函数原型33## 函数原型
34 34 
35-```cpp35+```c
36__aicore__ inline int64_t asc_get_sub_block_num()36__aicore__ inline int64_t asc_get_sub_block_num()
37```37```
38 38 
@@ -58,8 +58,78 @@ PIPE_S
58 58 
59## 约束说明59## 约束说明
60 60 
61-61+- 本接口读取的subblock维度保存在只读特殊寄存器中,由系统控制器(System Controller,SC)在kernel启动前配置,kernel运行期间不可修改,连续两次调用返回值相同。
62+- 返回值反映当前AI Core的subblock数量,仅在[分离模式](../../../../../guide/programming_guide/advanced_programming/hardware_implementation/basic_architecture.md)下有意义。非分离模式下返回值的语义以SC配置为准。
62 63 
63## 调用示例64## 调用示例
64 65 
65-66+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
67+ 
68+<!-- npu="950" id8 -->
69+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
70+ 
71+```bash
72+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
73+```
74+<!-- end id8 -->
75+ 
76+```cpp
77+#include <cstdint>
78+#include <iomanip>
79+#include <iostream>
80+#include <vector>
81+#include "c_api/asc_simd.h"
82+#include "acl/acl.h"
83+ 
84+namespace {
85+constexpr uint32_t ELEMENTS = 16;
86+constexpr uint32_t BYTES = ELEMENTS * sizeof(uint64_t);
87+ 
88+void print_data(const char* label, const std::vector<uint64_t>& data)
89+{
90+ std::cout << label << ":";
91+ for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << data[i];
92+ std::cout << " ..." << std::endl;
93+}
94+ 
95+void print_hex(const char* label, const std::vector<uint64_t>& data)
96+{
97+ std::cout << label << ":" << std::hex;
98+ for (uint32_t i = 0; i < 4; ++i) std::cout << " 0x" << data[i];
99+ std::cout << std::dec << std::endl;
100+}
101+ 
102+__global__ __vector__ void asc_get_sub_block_num_kernel(__gm__ uint64_t* output)
103+{
104+ asc_init();
105+ const uint64_t sub_id = static_cast<uint64_t>(asc_get_sub_block_id());
106+ const uint64_t sub_num = static_cast<uint64_t>(asc_get_sub_block_num());
107+ output[0] = sub_num;
108+ output[1] = sub_id;
109+ output[2] = sub_num;
110+ asc_dcci_single(output);
111+}
112+ 
113+} // namespace
114+ 
115+int main()
116+{
117+ std::vector<uint64_t> output(ELEMENTS, ~0ULL);
118+ aclInit(nullptr);
119+ aclrtSetDevice(0);
120+ uint64_t* output_device = nullptr;
121+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
122+ aclrtMemcpy(output_device, BYTES, output.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
123+ asc_get_sub_block_num_kernel<<<1, 0>>>(output_device);
124+ aclrtSynchronizeDevice();
125+ aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
126+ print_data("Observed", output);
127+ print_hex("Observed hex", output);
128+ const bool passed = output[2] > 0 && output[1] < output[2];
129+ std::cout << (passed ? "[Success] asc_get_sub_block_num passed." : "[Failed] asc_get_sub_block_num failed.") << std::endl;
130+ aclrtFree(output_device);
131+ aclrtResetDevice(0);
132+ aclFinalize();
133+ return passed ? 0 : 1;
134+}
135+```
@@ -28,11 +28,11 @@
28 28 
29头文件路径为:`"c_api/utils/sys_var.h"`29头文件路径为:`"c_api/utils/sys_var.h"`
30 30 
31-取系统虚拟基地址。31+取系统虚拟基地址特殊寄存器的值并按`int64_t`返回,该寄存器为只读特殊寄存器,存储当前核上下文的系统虚拟基地址,由系统在kernel启动前按核配置完成系统虚拟基地址用于在kernel内将相对偏移量转换为系统虚拟地址,常用于地址推导、跨核地址对齐等需要系统虚拟基地址参与的场景。
32 32 
33## 函数原型33## 函数原型
34 34 
35-```cpp35+```c
36 __aicore__ inline int64_t asc_get_sys_virtual_base()36 __aicore__ inline int64_t asc_get_sys_virtual_base()
37```37```
38 38 
@@ -50,11 +50,75 @@ PIPE_S
50 50 
51## 约束说明51## 约束说明
52 52 
53-53+- 本接口为只读查询接口,仅读取系统虚拟基地址寄存器当前值,不修改任何寄存器或存储状态,可在AIC与AIV上下文中调用。
54+- 本接口不支持纯SIMT编译模式。
54 55 
55## 调用示例56## 调用示例
56 57 
58+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
59+ 
60+<!-- npu="950" id8 -->
61+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
62+ 
63+```bash
64+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
65+```
66+<!-- end id8 -->
67+ 
57```cpp68```cpp
58-int64_t sys_virtual_base = asc_get_sys_virtual_base();69+#include <cstdint>
59-printf("sys virtual base is %x", sys_virtual_base);// 需用%x将其打印成十六进制的数70+#include <iomanip>
60-```71+#include <iostream>
72+#include <vector>
73+#include "c_api/asc_simd.h"
74+#include "acl/acl.h"
75+ 
76+namespace {
77+constexpr uint32_t ELEMENTS = 16;
78+constexpr uint32_t BYTES = ELEMENTS * sizeof(uint64_t);
79+ 
80+void print_data(const char* label, const std::vector<uint64_t>& data)
81+{
82+ std::cout << label << ":";
83+ for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << data[i];
84+ std::cout << " ..." << std::endl;
85+}
86+ 
87+void print_hex(const char* label, const std::vector<uint64_t>& data)
88+{
89+ std::cout << label << ":" << std::hex;
90+ for (uint32_t i = 0; i < 4; ++i) std::cout << " 0x" << data[i];
91+ std::cout << std::dec << std::endl;
92+}
93+ 
94+__global__ __vector__ void asc_get_sys_virtual_base_kernel(__gm__ uint64_t* output)
95+{
96+ asc_init();
97+ output[0] = static_cast<uint64_t>(asc_get_sys_virtual_base());
98+ output[1] = static_cast<uint64_t>(asc_get_sys_virtual_base());
99+ asc_dcci_single(output);
100+}
101+ 
102+} // namespace
103+ 
104+int main()
105+{
106+ std::vector<uint64_t> output(ELEMENTS, ~0ULL);
107+ aclInit(nullptr);
108+ aclrtSetDevice(0);
109+ uint64_t* output_device = nullptr;
110+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
111+ aclrtMemcpy(output_device, BYTES, output.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
112+ asc_get_sys_virtual_base_kernel<<<1, 0>>>(output_device);
113+ aclrtSynchronizeDevice();
114+ aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
115+ print_data("Observed", output);
116+ print_hex("Observed hex", output);
117+ const bool passed = output[0] == output[1];
118+ std::cout << (passed ? "[Success] asc_get_sys_virtual_base passed." : "[Failed] asc_get_sys_virtual_base failed.") << std::endl;
119+ aclrtFree(output_device);
120+ aclrtResetDevice(0);
121+ aclFinalize();
122+ return passed ? 0 : 1;
123+}
124+```
@@ -0,0 +1,141 @@
1+# asc_get_system_clock
M
Mmunanhw9 天前

新增文件要同步修改索引文件

likedislike
2+ 
3+## 产品支持情况
4+ 
5+<!-- npu="950" id1 -->
6+- Ascend 950PR/Ascend 950DT:支持
7+<!-- end id1 -->
8+<!-- npu="A3" id2 -->
9+- Atlas A3 训练系列产品/Atlas A3 推理系列产品:支持
10+<!-- end id2 -->
11+<!-- npu="910b" id3 -->
12+- Atlas A2 训练系列产品/Atlas A2 推理系列产品:支持
13+<!-- end id3 -->
14+<!-- npu="310b" id4 -->
15+- Atlas 200I/500 A2 推理产品:不支持
16+<!-- end id4 -->
17+<!-- npu="310p" id5 -->
18+- Atlas 推理系列产品AI Core:不支持
19+<!-- end id5 -->
20+<!-- npu="310p" id6 -->
21+- Atlas 推理系列产品Vector Core:不支持
22+<!-- end id6 -->
23+<!-- npu="910" id7 -->
24+- Atlas 训练系列产品:不支持
25+<!-- end id7 -->
26+ 
27+## 功能说明
28+ 
29+读取当前系统时钟计数器,返回`int64_t`类型的当前时间值(单位:微秒)。该接口在底层读取系统cycle计数器后,按当前芯片的时钟频率进行换算,直接返回微秒级时间值,无需调用方自行换算。
30+ 
31+<!-- npu="950" id8 -->
32+- 针对Ascend 950PR/Ascend 950DT,时钟频率为1GHz,换算系数为1000,即`clock = cycle / 1000`
33+<!-- end id8 -->
34+ 
35+<!-- npu="A3" id9 -->
36+- 针对Atlas A3 训练系列产品/Atlas A3 推理系列产品,时钟频率为50MHz,换算系数为50,即`clock = cycle / 50`
37+<!-- end id9 -->
38+ 
39+<!-- npu="910b" id10 -->
40+- 针对Atlas A2 训练系列产品/Atlas A2 推理系列产品,时钟频率为50MHz,换算系数为50,即`clock = cycle / 50`
41+<!-- end id10 -->
42+ 
43+本接口与[asc_get_system_cycle](asc_get_system_cycle.md)的区别在于:`asc_get_system_cycle`返回原始cycle计数值,需要调用方根据芯片频率自行换算为时间;`asc_get_system_clock`在接口内部完成换算,直接返回微秒级时间值。本接口在AIC与AIV上均可调用,返回值含义一致。
44+ 
45+## 函数原型
46+ 
47+```c
48+__aicore__ inline int64_t asc_get_system_clock()
49+```
50+ 
51+## 参数说明
52+ 
53+
54+ 
55+## 返回值说明
56+ 
57+返回`int64_t`类型的当前系统时间值,单位为微秒。该值由系统cycle计数器读数除以芯片对应的换算系数得到,随系统时钟持续递增。
58+ 
59+## 流水类型
60+ 
61+PIPE_S
62+ 
63+## 约束说明
64+ 
65+- 本接口为只读查询接口,读取的是系统时钟计数器并按芯片频率换算,不修改任何寄存器或存储状态,可在AIC与AIV上下文中调用。
66+- 返回值随系统时钟持续递增,同一kernel内多次调用返回值单调不减。
67+- 本接口与`asc_get_system_cycle`读取的是同一个底层计数器,区别仅在于本接口在内部完成了频率换算,返回值为微秒,而`asc_get_system_cycle`返回原始cycle值。
68+ 
69+## 调用示例
70+ 
71+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
72+ 
73+<!-- npu="950" id11 -->
74+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
75+ 
76+```bash
77+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
78+```
79+<!-- end id11 -->
80+ 
81+```cpp
82+#include <cstdint>
83+#include <iomanip>
84+#include <iostream>
85+#include <vector>
86+#include "c_api/asc_simd.h"
87+#include "acl/acl.h"
88+ 
89+namespace {
90+constexpr uint32_t ELEMENTS = 16;
91+constexpr uint32_t BYTES = ELEMENTS * sizeof(uint64_t);
92+ 
93+void print_data(const char* label, const std::vector<uint64_t>& data)
94+{
95+ std::cout << label << ":";
96+ for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << data[i];
97+ std::cout << " ..." << std::endl;
98+}
99+ 
100+void print_hex(const char* label, const std::vector<uint64_t>& data)
101+{
102+ std::cout << label << ":" << std::hex;
103+ for (uint32_t i = 0; i < 4; ++i) std::cout << " 0x" << data[i];
104+ std::cout << std::dec << std::endl;
105+}
106+ 
107+__global__ __vector__ void asc_get_system_clock_kernel(__gm__ uint64_t* output)
108+{
109+ asc_init();
110+ const uint64_t before = static_cast<uint64_t>(asc_get_system_clock());
111+ for (volatile uint32_t i = 0; i < 256; ++i) {}
112+ const uint64_t after = static_cast<uint64_t>(asc_get_system_clock());
113+ output[0] = before;
114+ output[1] = after;
115+ output[2] = after - before;
116+ asc_dcci_single(output);
117+}
118+ 
119+} // namespace
120+ 
121+int main()
122+{
123+ std::vector<uint64_t> output(ELEMENTS, ~0ULL);
124+ aclInit(nullptr);
125+ aclrtSetDevice(0);
126+ uint64_t* output_device = nullptr;
127+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
128+ aclrtMemcpy(output_device, BYTES, output.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
129+ asc_get_system_clock_kernel<<<1, 0>>>(output_device);
130+ aclrtSynchronizeDevice();
131+ aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
132+ print_data("Observed", output);
133+ print_hex("Observed hex", output);
134+ const bool passed = output[1] > output[0] && output[2] > 0;
135+ std::cout << (passed ? "[Success] asc_get_system_clock passed." : "[Failed] asc_get_system_clock failed.") << std::endl;
136+ aclrtFree(output_device);
137+ aclrtResetDevice(0);
138+ aclFinalize();
139+ return passed ? 0 : 1;
140+}
141+```
@@ -30,6 +30,9 @@
30 30 
31读取当前系统cycle计数器,返回int64_t类型的当前计数值。该计数器为只读系统计数器,宽度64bit,复位值为0,随系统时钟持续递增,反映硬件运行以来的累计cycle数,常用于性能统计、耗时测量与执行序的时序判断。本接口在AIC与AIV上均可调用,返回值含义一致。31读取当前系统cycle计数器,返回int64_t类型的当前计数值。该计数器为只读系统计数器,宽度64bit,复位值为0,随系统时钟持续递增,反映硬件运行以来的累计cycle数,常用于性能统计、耗时测量与执行序的时序判断。本接口在AIC与AIV上均可调用,返回值含义一致。
L
LLycheeeee14 天前

可以推荐使用新接口。

likedislike
32 32 
33+> [!NOTE]说明
34+> 如需直接获取微秒级时间值(接口内部完成频率换算),推荐使用[asc_get_system_clock](asc_get_system_clock.md)。
35+ 
33<!-- npu="950" id8 -->36<!-- npu="950" id8 -->
34- 针对Ascend 950PR/Ascend 950DT,若换算成时间需要按照1GHz的频率,时间单位为us,换算公式为:time = (cycle数/1000)us。37- 针对Ascend 950PR/Ascend 950DT,若换算成时间需要按照1GHz的频率,时间单位为us,换算公式为:time = (cycle数/1000)us。
35<!-- end id8 -->38<!-- end id8 -->
@@ -74,7 +77,7 @@ PIPE_S
74以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:77以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
75 78 
76```bash79```bash
77-bisheng example.asc -o main --npu-arch=dav-3510; ./main80+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
78```81```
79<!-- end id11 -->82<!-- end id11 -->
80 83 
@@ -28,11 +28,11 @@
28 28 
29头文件路径为:`"c_api/utils/sys_var.h"`29头文件路径为:`"c_api/utils/sys_var.h"`
30 30 
31-获取Tensor位宽VL(Vector Length)的大小31+获取当前芯片的矢量寄存器位宽VL([Vector Length](../../defs/type/data_type_definition.md),即矢量计算单元单次Repeat可处理元素个数VL是Reg矢量计算循环每轮消耗的元素数量,常用于推导mask长度、循环次数、单次搬运元素数等参数。本接口在AIC与AIV上均可调用,返回值含义一致。
32 32 
33## 函数原型33## 函数原型
34 34 
35-```cpp35+```c
36__aicore__ inline constexpr int64_t asc_get_vf_len()36__aicore__ inline constexpr int64_t asc_get_vf_len()
37```37```
38 38 
@@ -50,10 +50,77 @@ PIPE_S
50 50 
51## 约束说明51## 约束说明
52 52 
53-53+- 本接口不触发硬件指令,不依赖任何特殊寄存器或前置setter接口配置,可在AIC与AIV上下文中调用。
54+- 返回值反映当前芯片的矢量寄存器位宽,同一kernel内多次调用返回值相同。
55+- 本接口不支持纯SIMT编译模式,在该模式下包含对应头文件会触发编译错误。
54 56 
55## 调用示例57## 调用示例
56 58 
57-```cpp59+将代码保存为`example.asc`后,可通过`bisheng`命令编译运行,其中`--npu-arch`参数需根据实际产品型号指定对应的NPU架构,具体产品与NPU架构的映射关系请参考[\_\_NPU\_ARCH\_\_](../../../../../guide/programming_guide/language_extension/simd_builtin_keywords.md#npu-arch)。
58-int64_t len = asc_get_vf_len();60+ 
61+<!-- npu="950" id8 -->
62+以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
63+ 
64+```bash
65+bisheng example.asc -o main --npu-arch=dav-3510 && ./main
66+```
67+<!-- end id8 -->
68+ 
69+```cpp
70+#include <cstdint>
71+#include <iomanip>
72+#include <iostream>
73+#include <vector>
74+#include "c_api/asc_simd.h"
75+#include "acl/acl.h"
76+ 
77+namespace {
78+constexpr uint32_t ELEMENTS = 16;
79+constexpr uint32_t BYTES = ELEMENTS * sizeof(uint64_t);
80+ 
81+void print_data(const char* label, const std::vector<uint64_t>& data)
82+{
83+ std::cout << label << ":";
84+ for (uint32_t i = 0; i < 8; ++i) std::cout << ' ' << data[i];
85+ std::cout << " ..." << std::endl;
86+}
87+ 
88+void print_hex(const char* label, const std::vector<uint64_t>& data)
89+{
90+ std::cout << label << ":" << std::hex;
91+ for (uint32_t i = 0; i < 4; ++i) std::cout << " 0x" << data[i];
92+ std::cout << std::dec << std::endl;
93+}
94+ 
95+__global__ __vector__ void asc_get_vf_len_kernel(__gm__ uint64_t* output)
96+{
97+ asc_init();
98+ constexpr uint64_t vf_bytes = asc_get_vf_len();
99+ output[0] = vf_bytes;
100+ output[1] = vf_bytes / sizeof(float);
101+ asc_dcci_single(output);
102+}
103+ 
104+} // namespace
105+ 
106+int main()
107+{
108+ std::vector<uint64_t> output(ELEMENTS, ~0ULL);
109+ aclInit(nullptr);
110+ aclrtSetDevice(0);
111+ uint64_t* output_device = nullptr;
112+ aclrtMalloc(reinterpret_cast<void**>(&output_device), BYTES, ACL_MEM_MALLOC_HUGE_FIRST);
113+ aclrtMemcpy(output_device, BYTES, output.data(), BYTES, ACL_MEMCPY_HOST_TO_DEVICE);
114+ asc_get_vf_len_kernel<<<1, 0>>>(output_device);
115+ aclrtSynchronizeDevice();
116+ aclrtMemcpy(output.data(), BYTES, output_device, BYTES, ACL_MEMCPY_DEVICE_TO_HOST);
117+ print_data("Observed", output);
118+ print_hex("Observed hex", output);
119+ const bool passed = output[0] == 256 && output[1] == 64;
120+ std::cout << (passed ? "[Success] asc_get_vf_len passed." : "[Failed] asc_get_vf_len failed.") << std::endl;
121+ aclrtFree(output_device);
122+ aclrtResetDevice(0);
123+ aclFinalize();
124+ return passed ? 0 : 1;
125+}
59```126```
@@ -12,6 +12,7 @@
12- **[asc_get_sub_block_num](asc_get_sub_block_num.md)**12- **[asc_get_sub_block_num](asc_get_sub_block_num.md)**
13- **[asc_get_sys_virtual_base](asc_get_sys_virtual_base.md)**13- **[asc_get_sys_virtual_base](asc_get_sys_virtual_base.md)**
14- **[asc_get_system_cycle](asc_get_system_cycle.md)**14- **[asc_get_system_cycle](asc_get_system_cycle.md)**
15+- **[asc_get_system_clock](asc_get_system_clock.md)**
15- **[asc_get_vf_len](asc_get_vf_len.md)**16- **[asc_get_vf_len](asc_get_vf_len.md)**
16- **[asc_set_ffts_base_addr](asc_set_ffts_base_addr.md)**17- **[asc_set_ffts_base_addr](asc_set_ffts_base_addr.md)**
17- **[asc_get_block_idx(废弃)](../../deprecated_interface/asc_get_block_idx_deprecated.md)**18- **[asc_get_block_idx(废弃)](../../deprecated_interface/asc_get_block_idx_deprecated.md)**
@@ -88,7 +88,6 @@ PIPE_V
88 88 
89## 调用示例89## 调用示例
90 90 
91- 
92```cpp91```cpp
93// 示例计算256个half的和,分2次迭代完成,计算结果为2组求和的值92// 示例计算256个half的和,分2次迭代完成,计算结果为2组求和的值
94constexpr uint32_t total_length = 256;93constexpr uint32_t total_length = 256;
@@ -31,7 +31,6 @@
31和asc_copy_gm2ub_align接口配合使用,设置连续搬运数据块左右两侧需要填补的数据值。31和asc_copy_gm2ub_align接口配合使用,设置连续搬运数据块左右两侧需要填补的数据值。
32 32 
33<!-- npu="950" id10 -->33<!-- npu="950" id10 -->
34- 
35对于Ascend 950PR/Ascend 950DT产品:34对于Ascend 950PR/Ascend 950DT产品:
36- 该接口仅对asc_copy_gm2ub_align接口有效。35- 该接口仅对asc_copy_gm2ub_align接口有效。
37- fp8_e8m0_t、fp8_e5m2_t、fp8_e4m3fn_t、fp4x2_e2m1_t、fp4x2_e1m2_t、hifloat8_t类型的数据需转换成int8_t类型后再调用本接口。36- fp8_e8m0_t、fp8_e5m2_t、fp8_e4m3fn_t、fp4x2_e2m1_t、fp4x2_e1m2_t、hifloat8_t类型的数据需转换成int8_t类型后再调用本接口。
@@ -81,7 +81,6 @@ PIPE_S
81将以下代码保存为`example.asc`后,执行对应的编译运行命令。81将以下代码保存为`example.asc`后,执行对应的编译运行命令。
82 82 
83<!-- npu="950" id8 -->83<!-- npu="950" id8 -->
84- 
85以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:84以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
86 85 
87```bash86```bash
@@ -68,7 +68,6 @@ PIPE_S
68将以下代码保存为`example.asc`后,执行对应的编译运行命令。68将以下代码保存为`example.asc`后,执行对应的编译运行命令。
69 69 
70<!-- npu="950" id8 -->70<!-- npu="950" id8 -->
71- 
72以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:71以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
73 72 
74```bash73```bash
@@ -80,7 +80,6 @@ PIPE_S
80将以下代码保存为`example.asc`后,执行对应的编译运行命令。80将以下代码保存为`example.asc`后,执行对应的编译运行命令。
81 81 
82<!-- npu="950" id8 -->82<!-- npu="950" id8 -->
83- 
84以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:83以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
85 84 
86```bash85```bash
@@ -75,7 +75,6 @@ PIPE_S
75将以下代码保存为`example.asc`后,执行对应的编译运行命令。75将以下代码保存为`example.asc`后,执行对应的编译运行命令。
76 76 
77<!-- npu="950" id8 -->77<!-- npu="950" id8 -->
78- 
79以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:78以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
80 79 
81```bash80```bash
@@ -75,7 +75,6 @@ PIPE_S
75将以下代码保存为`example.asc`后,执行对应的编译运行命令。75将以下代码保存为`example.asc`后,执行对应的编译运行命令。
76 76 
77<!-- npu="950" id8 -->77<!-- npu="950" id8 -->
78- 
79以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:78以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
80 79 
81```bash80```bash
@@ -75,7 +75,6 @@ PIPE_S
75将以下代码保存为`example.asc`后,执行对应的编译运行命令。75将以下代码保存为`example.asc`后,执行对应的编译运行命令。
76 76 
77<!-- npu="950" id8 -->77<!-- npu="950" id8 -->
78- 
79以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:78以Ascend 950PR/Ascend 950DT产品(对应NPU架构为`dav-3510`)为例,编译运行命令如下:
80 79 
81```bash80```bash