Star704
1007
代码
代码
Issues82
Pull Requests181
Actions
项目讨论
Wiki
分析
Star704
1007
已合并
add reduce sample #1575
wang-zijiu创建于 4月13日
add reduce sample #1575
已合并
wang-zijiu创建于 4月13日
从wang-zijiu/asc-devkit_1889add_reg_br合入到cann/asc-devkitmaster
共 50 个文件变更+2238-107
Mexamples/01_simd_cpp_api/02_features/03_basic_api/00_data_movement/ld_st_reg_align/CMakeLists.txt+3-16
@@ -11,13 +11,8 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14-if(DEFINED RUN_MODE)14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15- set(CMAKE_ASC_RUN_MODE ${RUN_MODE} CACHE STRING "Alias of RUN_MODE")15+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
16-endif()
17- 
18-if(DEFINED NPU_ARCH)
19- set(CMAKE_ASC_ARCHITECTURES ${NPU_ARCH} CACHE STRING "Alias of NPU_ARCH")
20-endif()
21 16 
22find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
23 18 
@@ -34,14 +29,6 @@ target_link_libraries(demo PRIVATE
34 m29 m
35)30)
36 31 
37-# ======================================================================================
38-# NPU 编译选项配置
39-#
40-# 说明：
41-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
42-# ======================================================================================
43-set(NPU_ARCH "dav-3510" CACHE STRING "NPU ARCH, e.g. dav-3510")
44- 
45message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}")32message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}")
46 33 
47if(NOT DEFINED SCENARIO_NUM)34if(NOT DEFINED SCENARIO_NUM)
@@ -57,5 +44,5 @@ target_compile_definitions(demo PRIVATE
57)44)
58 45 
59target_compile_options(demo PRIVATE46target_compile_options(demo PRIVATE
60- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${NPU_ARCH}>47+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
61)48)
Mexamples/01_simd_cpp_api/02_features/03_basic_api/00_data_movement/ld_st_reg_align/README.md+6-6
@@ -114,17 +114,17 @@
114 ```bash114 ```bash
115 SCENARIO=1115 SCENARIO=1
116 mkdir -p build && cd build; # 创建并进入build目录116 mkdir -p build && cd build; # 创建并进入build目录
117- cmake -DNPU_ARCH=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # 编译工程（默认npu模式）117+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # 编译工程（默认npu模式）
118 ./demo # 执行编译生成的可执行程序，执行样例118 ./demo # 执行编译生成的可执行程序，执行样例
119 ```119 ```
120 120 
121- 使用 CPU调试 或 NPU仿真 模式时，添加 `-DRUN_MODE=cpu` 或 `-DRUN_MODE=sim` 参数即可。121+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
122 122 
123 示例如下：123 示例如下：
124 ```bash124 ```bash
125 SCENARIO=1125 SCENARIO=1
126- cmake -DRUN_MODE=cpu -DNPU_ARCH=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # cpu调试模式126+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # cpu调试模式
127- cmake -DRUN_MODE=sim -DNPU_ARCH=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # NPU仿真模式127+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # NPU仿真模式
128 ```128 ```
129 129 
130 > **注意：** 切换编译模式或场景前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。130 > **注意：** 切换编译模式或场景前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
@@ -133,8 +133,8 @@
133 133 
134| 选项　　　　　 | 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　 |134| 选项　　　　　 | 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　 |
135| ----------------| -----------------------------| --------------------------------------------------------------------------------------|135| ----------------| -----------------------------| --------------------------------------------------------------------------------------|
136-| `RUN_MODE`　　 | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　　　　　　　　　　　　　　　　　　 |136+| `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　　　　　　　　　　　　　　　　　　 |
137-| `NPU_ARCH`　　 | `dav-3510`　　　　　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT　　　　　　　　　　　　　　　　　　|137+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
138| `SCENARIO_NUM` | `1`、`2`、`3`、`4`　　　　　| 场景编号：1=连续搬运，2=连续搬运(postUpdate)，3=非连续搬运，4=非连续搬运(postUpdate) |138| `SCENARIO_NUM` | `1`、`2`、`3`、`4`　　　　　| 场景编号：1=连续搬运，2=连续搬运(postUpdate)，3=非连续搬运，4=非连续搬运(postUpdate) |
139 139 
140- 执行结果140- 执行结果
Mexamples/01_simd_cpp_api/02_features/03_basic_api/00_data_movement/ld_st_reg_align/ld_st_reg_align.asc+1-0
@@ -154,6 +154,7 @@ __global__ __vector__ void ld_st_reg_align_kernel(GM_ADDR x, GM_ADDR y)
154 Kernel<float> kernel;154 Kernel<float> kernel;
155 kernel.Init(x, y, 1024);155 kernel.Init(x, y, 1024);
156 kernel.Process();156 kernel.Process();
157+ AscendC::PipeBarrier<PIPE_ALL>();
157}158}
158 159 
159std::vector<float> ld_st_reg_align(std::vector<float>& input)160std::vector<float> ld_st_reg_align(std::vector<float>& input)
Mexamples/01_simd_cpp_api/02_features/03_basic_api/00_data_movement/ld_st_reg_mask/CMakeLists.txt+3-15
@@ -11,13 +11,8 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14-if(DEFINED RUN_MODE)14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15- set(CMAKE_ASC_RUN_MODE ${RUN_MODE} CACHE STRING "Alias of RUN_MODE")15+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
16-endif()
17- 
18-if(DEFINED NPU_ARCH)
19- set(CMAKE_ASC_ARCHITECTURES ${NPU_ARCH} CACHE STRING "Alias of NPU_ARCH")
20-endif()
21 16 
22find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
23 18 
@@ -27,13 +22,6 @@ add_executable(demo
27 ld_st_reg_mask.asc22 ld_st_reg_mask.asc
28)23)
29 24 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36-set(NPU_ARCH "dav-3510" CACHE STRING "NPU ARCH, e.g. dav-3510")
37target_compile_options(demo PRIVATE25target_compile_options(demo PRIVATE
38- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${NPU_ARCH}>26+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
39)27)
Mexamples/01_simd_cpp_api/02_features/03_basic_api/00_data_movement/ld_st_reg_mask/README.md+8-8
@@ -67,27 +67,27 @@
67- 样例执行67- 样例执行
68 ```bash68 ```bash
69 mkdir -p build && cd build; # 创建并进入build目录69 mkdir -p build && cd build; # 创建并进入build目录
70- cmake -DNPU_ARCH=dav-3510 ..;make -j; # 编译工程（默认npu模式）70+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程（默认npu模式）
71 python3 ../scripts/gen_data.py # 生成测试输入数据71 python3 ../scripts/gen_data.py # 生成测试输入数据
72 ./demo # 执行编译生成的可执行程序，执行样例72 ./demo # 执行编译生成的可执行程序，执行样例
73 ```73 ```
74 74 
75- 使用 CPU调试 或 NPU仿真 模式时，添加 `-DRUN_MODE=cpu` 或 `-DRUN_MODE=sim` 参数即可。75+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
76 76 
77 示例如下：77 示例如下：
78 ```bash78 ```bash
79- cmake -DRUN_MODE=cpu -DNPU_ARCH=dav-3510 ..;make -j; # cpu调试模式79+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
80- cmake -DRUN_MODE=sim -DNPU_ARCH=dav-3510 ..;make -j; # NPU仿真模式80+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
81 ```81 ```
82 82 
83 > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。83 > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
84 84 
85- 编译选项说明85- 编译选项说明
86 86 
87- | 选项 | 可选值 | 说明 |87+| 选项　　　　　　　　　　　| 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
88- |------|--------|------|88+| ---------------------------| -----------------------------| ---------------------------------------------------|
89- | `RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |89+| `CMAKE_ASC_RUN_MODE`　　　| `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
90- | `NPU_ARCH` | `dav-3510` | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |90+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510`　　　　　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
91 91 
92- 执行结果92- 执行结果
93 93 
Mexamples/01_simd_cpp_api/02_features/03_basic_api/00_data_movement/ld_st_reg_mask/ld_st_reg_mask.asc+1-0
@@ -82,6 +82,7 @@ __global__ __vector__ void ld_st_reg_mask(GM_ADDR x, GM_ADDR y)
82 Kernel<uint8_t> kernel;82 Kernel<uint8_t> kernel;
83 kernel.Init(x, y, 1024);83 kernel.Init(x, y, 1024);
84 kernel.Process();84 kernel.Process();
85+ AscendC::PipeBarrier<PIPE_ALL>();
85}86}
86 87 
87static bool CompareResult(const void* outputData, uint32_t outSize)88static bool CompareResult(const void* outputData, uint32_t outSize)
Mexamples/01_simd_cpp_api/02_features/03_basic_api/00_data_movement/ld_st_reg_unalign/CMakeLists.txt+3-15
@@ -11,13 +11,8 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14-if(DEFINED RUN_MODE)14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15- set(CMAKE_ASC_RUN_MODE ${RUN_MODE} CACHE STRING "Alias of RUN_MODE")15+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
16-endif()
17- 
18-if(DEFINED NPU_ARCH)
19- set(CMAKE_ASC_ARCHITECTURES ${NPU_ARCH} CACHE STRING "Alias of NPU_ARCH")
20-endif()
21 16 
22find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
23 18 
@@ -34,13 +29,6 @@ target_link_libraries(demo PRIVATE
34 m29 m
35)30)
36 31 
37-# ======================================================================================
38-# NPU 编译选项配置
39-#
40-# 说明：
41-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
42-# ======================================================================================
43-set(NPU_ARCH "dav-3510" CACHE STRING "NPU ARCH, e.g. dav-3510")
44target_compile_options(demo PRIVATE32target_compile_options(demo PRIVATE
45- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${NPU_ARCH}>33+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
46)34)
Mexamples/01_simd_cpp_api/02_features/03_basic_api/00_data_movement/ld_st_reg_unalign/README.md+6-6
@@ -57,16 +57,16 @@
57- 样例执行57- 样例执行
58 ```bash58 ```bash
59 mkdir -p build && cd build; # 创建并进入build目录59 mkdir -p build && cd build; # 创建并进入build目录
60- cmake -DNPU_ARCH=dav-3510 ..;make -j; # 编译工程60+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程
61 ./demo # 执行编译生成的可执行程序，执行样例61 ./demo # 执行编译生成的可执行程序，执行样例
62 ```62 ```
63 63 
64- 使用 CPU调试 或 NPU仿真 模式时，添加 `-DRUN_MODE=cpu` 或 `-DRUN_MODE=sim` 参数即可。64+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
65 65 
66 示例如下：66 示例如下：
67 ```bash67 ```bash
68- cmake -DRUN_MODE=cpu -DNPU_ARCH=dav-3510 ..;make -j; # cpu调试模式68+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
69- cmake -DRUN_MODE=sim -DNPU_ARCH=dav-3510 ..;make -j; # NPU仿真模式69+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
70 ```70 ```
71 71 
72 > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。72 > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
@@ -75,8 +75,8 @@
75 75 
76| 选项　　　 | 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|76| 选项　　　 | 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
77| ------------| -----------------------------| ---------------------------------------------------|77| ------------| -----------------------------| ---------------------------------------------------|
78-| `RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|78+| `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
79-| `NPU_ARCH` | `dav-3510`　　　　　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |79+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
80 80 
81- 执行结果81- 执行结果
82 82 
Mexamples/01_simd_cpp_api/02_features/03_basic_api/00_data_movement/ld_st_reg_unalign/ld_st_reg_unalign.asc+1-0
@@ -86,6 +86,7 @@ __global__ __vector__ void ld_st_reg_unalign_kernel(GM_ADDR x, GM_ADDR y)
86 Kernel<float> kernel;86 Kernel<float> kernel;
87 kernel.Init(x, y, 1024);87 kernel.Init(x, y, 1024);
88 kernel.Process();88 kernel.Process();
89+ AscendC::PipeBarrier<PIPE_ALL>();
89}90}
90 91 
91std::vector<float> ld_st_reg_unalign(std::vector<float>& input)92std::vector<float> ld_st_reg_unalign(std::vector<float>& input)
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/README.md+20-0
@@ -0,0 +1,20 @@
1+# 向量计算类api样例介绍
2+ 
3+## 概述
4+ 
5+本路径下包含了与向量计算相关的多个API的样例。每个样例均基于Ascend C的<<<>>>直调方法，支持main函数和kernel函数在同一个cpp文件中实现。
6+ 
7+## 算子开发样例
8+ 
9+| 目录名称　　　　　　　　　　　 | 功能描述　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　　|
10+| --------------------------------| ---------------------------------------------------------------------------------------------|
11+| [abs](./abs)　　　　　　　　　 | 本样例基于RegBase编程范式实现Abs运算，Relu/Exp/Sqrt/Ln/Log/Log2/Log10/Neg接口皆可参考该样例 |
12+| [mul](./mul)　　　　　　　　　 | 本样例基于RegBase编程范式实现Mul运算，Add/Sub/Div/Max/Min/Prelu接口皆可参考该样例　　　　　 |
13+| [muls](./muls)　　　　　　　　 | 本样例基于RegBase编程范式实现Muls运算，Adds/Maxs/Mins/LeakyRelu接口皆可参考该样例　　　　　 |
14+| [reduce](./reduce)　　　　　　 | 本样例基于RegBase编程范式实现Reduce运算（SUM模式），Reduce接口支持SUM/MAX/MIN归约模式　　　 |
15+| [reduce_pair](./reduce_pair)　 | 本样例基于RegBase编程范式实现PairReduceElem运算（SUM模式），对相邻奇偶元素对进行归约求和　　|
16+| [reduce_block](./reduce_block) | 本样例基于RegBase编程范式实现ReduceDataBlock运算（SUM模式），对每个DataBlock(32B)内元素归约 |
17+| [cast](./cast)　　　　　　　　 | 本样例基于RegBase编程范式实现Cast运算，用于RegTensor数据类型转换（位宽大转小/小转大）　　　 |
18+| [truncate](./truncate)　　　　 | 本样例基于RegBase编程范式实现Truncate运算，将浮点数截断到整数位　　　　　　　　　　　　　　 |
19+| [arange](./arange)　　　　　　 | 本样例基于RegBase编程范式实现Arange运算，以标量值为起始生成递增/递减索引序列　　　　　　　　|
20+| [duplicate](./duplicate)　　　 | 本样例基于RegBase编程范式实现Duplicate运算（标量填充），将标量值填充到向量的每个位置　　　　|
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/abs/abs.asc+6-1
@@ -1,5 +1,5 @@
1/**1/**
2-* Copyright (c) 2025 Huawei Technologies Co., Ltd.2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3* This program is free software, you can redistribute it and/or modify it under the terms and conditions of3* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4* CANN Open Software License Agreement Version 2.0 (the "License").4* CANN Open Software License Agreement Version 2.0 (the "License").
5* Please refer to the License for details. You may not use this file except in compliance with the License.5* Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -9,6 +9,11 @@
9*/9*/
10 10 
11 11 
12+/* !
13+ * \file abs.asc
14+ * \brief 演示使用Abs API实现向量绝对值计算功能，对输入向量逐元素取绝对值并输出结果
15+ */
16+ 
12#include "acl/acl.h"17#include "acl/acl.h"
13#include "kernel_operator.h"18#include "kernel_operator.h"
14#include "data_utils.h"19#include "data_utils.h"
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/abs/data_utils.h+1-1
@@ -11,7 +11,7 @@
11 11 
12/* !12/* !
13 * \file data_utils.h13 * \file data_utils.h
14- * \brief14+ * \brief Abs样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15 */15 */
16 16 
17#ifndef DATA_UTILS_H17#ifndef DATA_UTILS_H
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/arange/CMakeLists.txt+35-0
@@ -0,0 +1,35 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+cmake_minimum_required(VERSION 3.16)
13+ 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
16+ 
17+find_package(ASC REQUIRED)
18+ 
19+project(kernel_samples LANGUAGES ASC CXX)
20+ 
21+add_executable(demo
22+ arange.asc
23+)
24+ 
25+target_link_libraries(demo PRIVATE
26+ tiling_api
27+ register
28+ platform
29+ m
30+)
31+ 
32+target_compile_options(demo PRIVATE
33+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
34+ $<$<COMPILE_LANGUAGE:ASC>:--cce-simd-vf-fusion=false>
35+)
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/arange/README.md+94-0
@@ -0,0 +1,94 @@
1+# arange样例
2+ 
3+## 概述
4+本样例基于RegBase编程范式实现Arange运算，主要调用Arange接口。
5+- Arange接口以传入的标量值为起始值，生成递增/递减的索引序列
6+- 本样例以递增模式为例，起始值为0
7+ 
8+## 支持的产品
9+- Ascend 950PR/Ascend 950DT
10+ 
11+## 目录结构介绍
12+```
13+├── arange
14+│ ├── scripts
15+│ │ │ ├── gen_data.py // 真值数据生成脚本
16+│ ├── CMakeLists.txt // 编译工程文件
17+│ ├── data_utils.h // 数据读入写出函数
18+│ ├── arange.asc // AscendC样例实现 & 调用样例
19+│ └── README.md // 样例介绍
20+```
21+ 
22+## 样例描述
23+- 样例功能：
24+ 生成从0开始的递增索引序列，向量shape为[1, 256]，数据类型为float。
25+ Arange以起始值为0生成递增索引：{0, 1, 2, 3, ..., 255}。
26+- 样例规格：
27+ <table>
28+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="3" align="center">AIV样例</td></tr>
29+ <tr><td rowspan="2" align="center">样例输出</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
30+ <tr><td align="center">y</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
R
Rruanyy4月15日

表格错位了

likedislike
31+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">arange</td></tr>
32+ </table>
33+- 样例实现：
34+ ArangeVF函数内调用Arange接口进行索引生成：
35+ - 使用Arange接口生成递增索引，每次repeat生成oneRepeatSize个索引
36+ - 每次repeat后更新起始值
37+ - 调用实现
38+ 使用内核调用符<<<>>>调用核函数。
R
Rruanyy4月15日

空行未生效

likedislike
39+ 
40+## 注意事项
41+- Arange接口不需要输入数据，只需指定起始标量值
42+- 每次repeat生成oneRepeatSize个递增索引，多次repeat时需要手动更新起始值
43+- 默认模板参数为递增模式（IndexOrder::INCREASE_ORDER）
44+ 
45+## 编译运行
46+在本样例根目录下执行如下步骤，编译并执行样例。
47+- 配置环境变量
48+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
49+ - 默认路径，root用户安装CANN软件包
50+ ```bash
51+ source /usr/local/Ascend/cann/set_env.sh
52+ ```
53+ 
54+ - 默认路径，非root用户安装CANN软件包
55+ ```bash
56+ source $HOME/Ascend/cann/set_env.sh
57+ ```
58+ 
59+ - 指定路径install_path，安装CANN软件包
60+ ```bash
61+ source ${install_path}/cann/set_env.sh
62+ ```
63+ 
64+- 样例执行
65+ ```bash
66+ mkdir -p build && cd build; # 创建并进入build目录
67+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程（默认npu模式）
68+ python3 ../scripts/gen_data.py # 生成测试真值数据
69+ ./demo # 执行编译生成的可执行程序，执行样例
70+ ```
71+ 
72+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
73+ 
74+ 示例如下：
75+ ```bash
76+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
77+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
78+ ```
79+ 
80+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
81+ 
82+- 编译选项说明
83+ 
84+| 选项　　　　　　　　　　　| 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
85+| ---------------------------| -----------------------------| ---------------------------------------------------|
86+| `CMAKE_ASC_RUN_MODE`　　　| `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
87+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510`　　　　　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
88+ 
89+- 执行结果
90+ 
91+ 执行结果如下，说明精度对比成功。
92+ ```bash
93+ test pass!
94+ ```
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/arange/arange.asc+149-0
@@ -0,0 +1,149 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file arange.asc
14+ * \brief 演示使用Arange API实现等差数列生成功能，根据起始值和步长生成连续的等差序列
15+ */
16+ 
17+#include "acl/acl.h"
18+#include "kernel_operator.h"
19+#include "data_utils.h"
20+ 
21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
25+template <typename T>
26+__simd_vf__ inline void ArangeVF(__ubuf__ T* dstAddr, T scalarValue, uint32_t oneRepeatSize)
27+{
28+ AscendC::Reg::RegTensor<T> dstReg;
29+ AscendC::Reg::MaskReg mask = AscendC::Reg::CreateMask<T>();
30+ // Arange: 以scalarValue为起始值，生成递增索引
31+ // 默认递增模式: {scalarValue, scalarValue+1, scalarValue+2, ...}
32+ AscendC::Reg::Arange(dstReg, scalarValue);
33+ AscendC::Reg::StoreAlign(dstAddr, dstReg, mask);
34+}
35+ 
36+template <typename T>
37+class Kernel {
38+public:
39+ __aicore__ inline void Init(GM_ADDR y, uint32_t totalLength)
40+ {
41+ this->totalLength = totalLength;
42+ this->yGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(y));
43+ }
44+ __aicore__ inline void Process()
45+ {
46+ AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator;
47+ AscendC::LocalTensor<T> yLocal = ubAllocator.Alloc<T, 256>();
48+ 
49+ // Compute: RegBase ArangeVF
50+ __ubuf__ T* yAddr = reinterpret_cast<__ubuf__ T*>(yLocal.GetPhyAddr());
51+ // 一次repeat计算的数据个数
52+ constexpr uint32_t oneRepeatSize = AscendC::GetVecLen() / sizeof(T);
53+ // 总迭代次数
54+ uint16_t repeatTimes = AscendC::CeilDivision(totalLength, oneRepeatSize);
55+ // 起始值为0，每次repeat后递增oneRepeatSize
56+ T startValue = static_cast<T>(0);
57+ for (uint16_t i = 0; i < repeatTimes; ++i) {
58+ ArangeVF(yAddr + i * oneRepeatSize, startValue, oneRepeatSize);
59+ startValue += static_cast<T>(oneRepeatSize);
60+ }
61+ 
62+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
63+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
64+ 
65+ // CopyOut: UB -> GM
66+ AscendC::DataCopy(yGm, yLocal, totalLength);
67+ }
68+ 
69+private:
70+ uint32_t totalLength;
71+ AscendC::GlobalTensor<T> yGm;
72+};
73+ 
74+__global__ __vector__ void arange(GM_ADDR y)
75+{
76+ AscendC::InitSocState();
77+ uint32_t totalLength = 256;
78+ Kernel<float> kernel;
79+ kernel.Init(y, totalLength);
80+ kernel.Process();
81+ AscendC::PipeBarrier<PIPE_ALL>();
82+}
83+ 
84+static bool CompareResult(const void* outputData, uint32_t outSize)
85+{
86+ void* goldenData;
87+ aclrtMallocHost((void**)(&goldenData), outSize);
88+ size_t goldenSize = outSize;
89+ bool ret = ReadFile("./output/golden.bin", goldenSize, goldenData, goldenSize);
90+ if (ret) {
91+ printf("ReadFile golden.bin success!\n");
92+ } else {
93+ printf("test failed!\n");
94+ return false;
95+ }
96+ constexpr float EPS = 1e-4;
97+ int64_t wrongNum = 0;
98+ 
99+ for (size_t i = 0; i < outSize / sizeof(float); i++) {
100+ float a = (reinterpret_cast<const float*>(outputData))[i];
101+ float b = (reinterpret_cast<const float*>(goldenData))[i];
102+ float ae = std::abs(a - b);
103+ float re = ae / std::abs(b);
104+ if (ae > EPS && re > EPS) {
105+ printf("CompareResult golden.bin failed at index %zu: output is %lf, golden is %lf\n", i, a, b);
106+ wrongNum++;
107+ }
108+ }
109+ aclrtFreeHost(goldenData);
110+ if (wrongNum != 0) {
111+ return false;
112+ } else {
113+ printf("CompareResult golden.bin success!\n");
114+ return true;
115+ }
116+}
117+ 
118+int32_t main(int32_t argc, char* argv[])
119+{
120+ uint32_t numBlocks = 1;
121+ aclInit(nullptr);
122+ int32_t deviceId = 0;
123+ aclrtSetDevice(deviceId);
124+ aclrtStream stream = nullptr;
125+ aclrtCreateStream(&stream);
126+ uint32_t totalLength = 256;
127+ size_t outputByteSize = static_cast<size_t>(1) * totalLength * sizeof(float);
128+ uint8_t *yHost;
129+ uint8_t *yDevice;
130+ aclrtMallocHost((void**)(&yHost), outputByteSize);
131+ aclrtMalloc((void**)&yDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
132+ arange<<<numBlocks, nullptr, stream>>>(yDevice);
133+ aclrtSynchronizeStream(stream);
134+ aclrtMemcpy(yHost, outputByteSize, yDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
135+ WriteFile("./output/output.bin", yHost, outputByteSize);
136+ bool goldenResult = true;
137+ goldenResult = CompareResult(yHost, outputByteSize);
138+ if (goldenResult) {
139+ printf("test pass!\n");
140+ } else {
141+ printf("test failed!\n");
142+ }
143+ aclrtFree(yDevice);
144+ aclrtFreeHost(yHost);
145+ aclrtDestroyStream(stream);
146+ aclrtResetDevice(deviceId);
147+ aclFinalize();
148+ return 0;
149+}
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/arange/data_utils.h+94-0
@@ -0,0 +1,94 @@
1+/**
2+* Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file data_utils.h
14+ * \brief Arange样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15+ */
16+ 
17+#ifndef DATA_UTILS_H
18+#define DATA_UTILS_H
19+#include <fcntl.h>
20+#include <sys/stat.h>
21+#include <unistd.h>
22+#include <fstream>
23+ 
24+#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
25+ 
26+bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize)
27+{
28+ struct stat sBuf;
29+ int fileStatus = stat(filePath.data(), &sBuf);
30+ if (fileStatus == -1) {
31+ ERROR_LOG("failed to get file");
32+ return false;
33+ }
34+ if (S_ISREG(sBuf.st_mode) == 0) {
35+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
36+ return false;
37+ }
38+ 
39+ std::ifstream file;
40+ file.open(filePath, std::ios::binary);
41+ if (!file.is_open()) {
42+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
43+ return false;
44+ }
45+ 
46+ std::filebuf *buf = file.rdbuf();
47+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
48+ if (size == 0) {
49+ ERROR_LOG("file size is 0");
50+ file.close();
51+ return false;
52+ }
53+ if (size > bufferSize) {
54+ ERROR_LOG("file size is larger than buffer size");
55+ file.close();
56+ return false;
57+ }
58+ buf->pubseekpos(0, std::ios::in);
59+ buf->sgetn(static_cast<char *>(buffer), size);
60+ fileSize = size;
61+ file.close();
62+ return true;
63+}
64+ 
65+/**
66+ * @brief Write data to file
67+ * @param [in] filePath: file path
68+ * @param [in] buffer: data to write to file
69+ * @param [in] size: size to write
70+ * @return write result
71+ */
72+bool WriteFile(const std::string &filePath, const void *buffer, size_t size)
73+{
74+ if (buffer == nullptr) {
75+ ERROR_LOG("Write file failed. buffer is nullptr");
76+ return false;
77+ }
78+ 
79+ int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE);
80+ if (fd < 0) {
81+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
82+ return false;
83+ }
84+ 
85+ size_t writeSize = write(fd, buffer, size);
86+ (void)close(fd);
87+ if (writeSize != size) {
88+ ERROR_LOG("Write file Failed.");
89+ return false;
90+ }
91+ 
92+ return true;
93+}
94+#endif // DATA_UTILS_H
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/arange/scripts/gen_data.py+28-0
@@ -0,0 +1,28 @@
1+#!/usr/bin/python3
2+# coding=utf-8
3+ 
4+# ----------------------------------------------------------------------------------------------------------
5+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
C
Cchangxianyu4月16日

2026，统一排查

likedislike
6+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
7+# CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
10+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR OR FITNESS FOR A PARTICULAR PURPOSE.
11+# See LICENSE in the root of the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------------------------------------
13+ 
14+ 
15+import os
16+import numpy as np
17+ 
18+def gen_golden_data_simple():
19+ total_length = 256
20+ data_type = np.float32
21+ # Arange: 从0开始生成递增索引
22+ # golden = {0, 1, 2, 3, ..., 255}
23+ golden = np.arange(0, total_length, 1).astype(data_type)
24+ os.makedirs("output", exist_ok=True)
25+ golden.tofile('./output/golden.bin')
26+ 
27+if __name__ == "__main__":
28+ gen_golden_data_simple()
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/cast/cast.asc+5-0
@@ -9,6 +9,11 @@
9*/9*/
10 10 
11 11 
12+/* !
13+ * \file cast.asc
14+ * \brief 演示使用Cast API实现数据类型转换功能，将输入向量从一种数据类型转换为另一种数据类型
15+ */
16+ 
12#include "acl/acl.h"17#include "acl/acl.h"
13#include "kernel_operator.h"18#include "kernel_operator.h"
14#include "data_utils.h"19#include "data_utils.h"
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/cast/data_utils.h+1-1
@@ -11,7 +11,7 @@
11 11 
12/* !12/* !
13 * \file data_utils.h13 * \file data_utils.h
14- * \brief14+ * \brief Cast样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15 */15 */
16 16 
17#ifndef DATA_UTILS_H17#ifndef DATA_UTILS_H
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/duplicate/CMakeLists.txt+35-0
@@ -0,0 +1,35 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+cmake_minimum_required(VERSION 3.16)
13+ 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
16+ 
17+find_package(ASC REQUIRED)
18+ 
19+project(kernel_samples LANGUAGES ASC CXX)
20+ 
21+add_executable(demo
22+ duplicate.asc
23+)
24+ 
25+target_link_libraries(demo PRIVATE
26+ tiling_api
27+ register
28+ platform
29+ m
30+)
31+ 
32+target_compile_options(demo PRIVATE
33+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
34+ $<$<COMPILE_LANGUAGE:ASC>:--cce-simd-vf-fusion=false>
35+)
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/duplicate/README.md+88-0
@@ -0,0 +1,88 @@
1+# duplicate样例
2+ 
3+## 概述
4+本样例基于RegBase编程范式实现Duplicate运算（标量填充模式），主要调用Duplicate接口。
5+- `Duplicate(dstReg, scalarValue, mask)` 接口将一个标量值复制多次并填充到向量中
6+- `Duplicate(dstReg, srcReg, mask)` 接口将源RegTensor的第一个元素复制多次并填充到向量中
7+ 
8+## 支持的产品
9+- Ascend 950PR/Ascend 950DT
10+ 
11+## 目录结构介绍
12+```
13+├── duplicate
14+│ ├── scripts
15+│ │ │ ├── gen_data.py // 真值数据生成脚本
16+│ ├── CMakeLists.txt // 编译工程文件
17+│ ├── data_utils.h // 数据读入写出函数
18+│ ├── duplicate.asc // AscendC样例实现 & 调用样例
19+│ └── README.md // 样例介绍
20+```
21+ 
22+## 样例描述
23+- 样例功能：
24+ 将标量值3.14填充到输出向量的每个位置，向量元素个数为256，数据类型为float。
25+- 样例规格：
26+ <table>
27+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="3" align="center">AIV样例</td></tr>
28+ <tr><td rowspan="2" align="center">样例输出</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
29+ <tr><td align="center">y</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
30+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">duplicate</td></tr>
31+ </table>
32+- 样例实现：
33+ DuplicateScalarVF函数内调用Duplicate接口进行标量填充：
34+ - 使用Duplicate接口将标量值填充到dstReg的每个元素
35+ - 使用StoreAlign将结果写回UB
36+ - 调用实现
37+ 使用内核调用符<<<>>>调用核函数。
38+ 
39+## 编译运行
40+在本样例根目录下执行如下步骤，编译并执行样例。
41+- 配置环境变量
42+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
43+ - 默认路径，root用户安装CANN软件包
44+ ```bash
45+ source /usr/local/Ascend/cann/set_env.sh
46+ ```
47+ 
48+ - 默认路径，非root用户安装CANN软件包
49+ ```bash
50+ source $HOME/Ascend/cann/set_env.sh
51+ ```
52+ 
53+ - 指定路径install_path，安装CANN软件包
54+ ```bash
55+ source ${install_path}/cann/set_env.sh
56+ ```
57+ 
58+- 样例执行
59+ ```bash
60+ mkdir -p build && cd build; # 创建并进入build目录
61+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程（默认npu模式）
62+ python3 ../scripts/gen_data.py # 生成测试真值数据
63+ ./demo # 执行编译生成的可执行程序，执行样例
64+ ```
65+ 
66+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
67+ 
68+ 示例如下：
69+ ```bash
70+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
71+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
72+ ```
73+ 
74+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
75+ 
76+- 编译选项说明
77+ 
78+| 选项　　　　　　　　　　　| 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
79+| ---------------------------| -----------------------------| ---------------------------------------------------|
80+| `CMAKE_ASC_RUN_MODE`　　　| `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
81+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510`　　　　　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
82+ 
83+- 执行结果
84+ 
85+ 执行结果如下，说明精度对比成功。
86+ ```bash
87+ test pass!
88+ ```
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/duplicate/data_utils.h+94-0
@@ -0,0 +1,94 @@
1+/**
2+* Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file data_utils.h
14+ * \brief Duplicate样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15+ */
16+ 
17+#ifndef DATA_UTILS_H
18+#define DATA_UTILS_H
19+#include <fcntl.h>
20+#include <sys/stat.h>
21+#include <unistd.h>
22+#include <fstream>
23+ 
24+#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
25+ 
26+bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize)
27+{
28+ struct stat sBuf;
29+ int fileStatus = stat(filePath.data(), &sBuf);
30+ if (fileStatus == -1) {
31+ ERROR_LOG("failed to get file");
32+ return false;
33+ }
34+ if (S_ISREG(sBuf.st_mode) == 0) {
35+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
36+ return false;
37+ }
38+ 
39+ std::ifstream file;
40+ file.open(filePath, std::ios::binary);
41+ if (!file.is_open()) {
42+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
43+ return false;
44+ }
45+ 
46+ std::filebuf *buf = file.rdbuf();
47+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
48+ if (size == 0) {
49+ ERROR_LOG("file size is 0");
50+ file.close();
51+ return false;
52+ }
53+ if (size > bufferSize) {
54+ ERROR_LOG("file size is larger than buffer size");
55+ file.close();
56+ return false;
57+ }
58+ buf->pubseekpos(0, std::ios::in);
59+ buf->sgetn(static_cast<char *>(buffer), size);
60+ fileSize = size;
61+ file.close();
62+ return true;
63+}
64+ 
65+/**
66+ * @brief Write data to file
67+ * @param [in] filePath: file path
68+ * @param [in] buffer: data to write to file
69+ * @param [in] size: size to write
70+ * @return write result
71+ */
72+bool WriteFile(const std::string &filePath, const void *buffer, size_t size)
73+{
74+ if (buffer == nullptr) {
75+ ERROR_LOG("Write file failed. buffer is nullptr");
76+ return false;
77+ }
78+ 
79+ int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE);
80+ if (fd < 0) {
81+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
82+ return false;
83+ }
84+ 
85+ size_t writeSize = write(fd, buffer, size);
86+ (void)close(fd);
87+ if (writeSize != size) {
88+ ERROR_LOG("Write file Failed.");
89+ return false;
90+ }
91+ 
92+ return true;
93+}
94+#endif // DATA_UTILS_H
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/duplicate/duplicate.asc+149-0
@@ -0,0 +1,149 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file duplicate.asc
14+ * \brief 演示使用Duplicate API实现数据填充功能，将一个标量值或立即数复制多次并填充到向量中
15+ */
16+ 
17+#include "acl/acl.h"
18+#include "kernel_operator.h"
19+#include "data_utils.h"
20+ 
21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
25+template <typename T>
26+__simd_vf__ inline void DuplicateScalarVF(__ubuf__ T* dstAddr, T scalarValue,
27+ uint32_t repeatTimes, uint32_t oneRepeatSize)
28+{
29+ AscendC::Reg::RegTensor<T> dstReg;
30+ AscendC::Reg::MaskReg mask = AscendC::Reg::CreateMask<T>();
31+ for (uint16_t i = 0; i < repeatTimes; ++i) {
32+ // Duplicate: 将标量值填充到dstReg的每个元素
33+ // dst[i] = scalarValue, for all i
34+ AscendC::Reg::Duplicate(dstReg, scalarValue, mask);
35+ AscendC::Reg::StoreAlign(dstAddr + i * oneRepeatSize, dstReg, mask);
36+ }
37+}
38+ 
39+template <typename T>
40+class Kernel {
41+public:
42+ __aicore__ inline void Init(GM_ADDR y, uint32_t totalLength)
43+ {
44+ this->totalLength = totalLength;
45+ this->yGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(y));
46+ }
47+ __aicore__ inline void Process()
48+ {
49+ AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator;
50+ AscendC::LocalTensor<T> yLocal = ubAllocator.Alloc<T, 256>();
51+ 
52+ // Compute: RegBase DuplicateScalarVF
53+ __ubuf__ T* yAddr = reinterpret_cast<__ubuf__ T*>(yLocal.GetPhyAddr());
54+ // 一次repeat计算的数据个数
55+ constexpr uint32_t oneRepeatSize = AscendC::GetVecLen() / sizeof(T);
56+ // 总迭代次数
57+ uint16_t repeatTimes = AscendC::CeilDivision(totalLength, oneRepeatSize);
58+ // 填充的标量值
59+ T fillValue = static_cast<T>(3.14);
60+ DuplicateScalarVF(yAddr, fillValue, repeatTimes, oneRepeatSize);
61+ 
62+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
63+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
64+ 
65+ // CopyOut: UB -> GM
66+ AscendC::DataCopy(yGm, yLocal, totalLength);
67+ }
68+ 
69+private:
70+ uint32_t totalLength;
71+ AscendC::GlobalTensor<T> yGm;
72+};
73+ 
74+__global__ __vector__ void duplicate(GM_ADDR y)
75+{
76+ AscendC::InitSocState();
77+ uint32_t totalLength = 256;
78+ Kernel<float> kernel;
79+ kernel.Init(y, totalLength);
80+ kernel.Process();
81+ AscendC::PipeBarrier<PIPE_ALL>();
82+}
83+ 
84+static bool CompareResult(const void* outputData, uint32_t outSize)
85+{
86+ void* goldenData;
87+ aclrtMallocHost((void**)(&goldenData), outSize);
88+ size_t goldenSize = outSize;
89+ bool ret = ReadFile("./output/golden.bin", goldenSize, goldenData, goldenSize);
90+ if (ret) {
91+ printf("ReadFile golden.bin success!\n");
92+ } else {
93+ printf("test failed!\n");
94+ return false;
95+ }
96+ constexpr float EPS = 1e-4;
97+ int64_t wrongNum = 0;
98+ 
99+ for (size_t i = 0; i < outSize / sizeof(float); i++) {
100+ float a = (reinterpret_cast<const float*>(outputData))[i];
101+ float b = (reinterpret_cast<const float*>(goldenData))[i];
102+ float ae = std::abs(a - b);
103+ float re = ae / std::abs(b);
104+ if (ae > EPS && re > EPS) {
105+ printf("CompareResult golden.bin failed at index %zu: output is %lf, golden is %lf\n", i, a, b);
106+ wrongNum++;
107+ }
108+ }
109+ aclrtFreeHost(goldenData);
110+ if (wrongNum != 0) {
111+ return false;
112+ } else {
113+ printf("CompareResult golden.bin success!\n");
114+ return true;
115+ }
116+}
117+ 
118+int32_t main(int32_t argc, char* argv[])
119+{
120+ uint32_t numBlocks = 1;
121+ aclInit(nullptr);
122+ int32_t deviceId = 0;
123+ aclrtSetDevice(deviceId);
124+ aclrtStream stream = nullptr;
125+ aclrtCreateStream(&stream);
126+ uint32_t totalLength = 256;
127+ size_t outputByteSize = static_cast<size_t>(1) * totalLength * sizeof(float);
128+ uint8_t *yHost;
129+ uint8_t *yDevice;
130+ aclrtMallocHost((void**)(&yHost), outputByteSize);
131+ aclrtMalloc((void**)&yDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
132+ duplicate<<<numBlocks, nullptr, stream>>>(yDevice);
133+ aclrtSynchronizeStream(stream);
134+ aclrtMemcpy(yHost, outputByteSize, yDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
135+ WriteFile("./output/output.bin", yHost, outputByteSize);
136+ bool goldenResult = true;
137+ goldenResult = CompareResult(yHost, outputByteSize);
138+ if (goldenResult) {
139+ printf("test pass!\n");
140+ } else {
141+ printf("test failed!\n");
142+ }
143+ aclrtFree(yDevice);
144+ aclrtFreeHost(yHost);
145+ aclrtDestroyStream(stream);
146+ aclrtResetDevice(deviceId);
147+ aclFinalize();
148+ return 0;
149+}
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/duplicate/scripts/gen_data.py+28-0
@@ -0,0 +1,28 @@
1+#!/usr/bin/python3
2+# coding=utf-8
3+ 
4+# ----------------------------------------------------------------------------------------------------------
5+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
7+# CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
10+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY OR OR FITNESS FOR A PARTICULAR PURPOSE.
11+# See LICENSE in the root of the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------------------------------------
13+ 
14+ 
15+import os
16+import numpy as np
17+ 
18+def gen_golden_data_simple():
19+ total_length = 256
20+ data_type = np.float32
21+ # DuplicateScalar: 将标量值3.14填充到所有位置
22+ fill_value = 3.14
23+ golden = np.full([1, total_length], fill_value, dtype=data_type)
24+ os.makedirs("output", exist_ok=True)
25+ golden.tofile('./output/golden.bin')
26+ 
27+if __name__ == "__main__":
28+ gen_golden_data_simple()
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/mul/README.md+5-5
@@ -57,7 +57,7 @@
57- 样例执行57- 样例执行
58 ```bash58 ```bash
59 mkdir -p build && cd build; # 创建并进入build目录59 mkdir -p build && cd build; # 创建并进入build目录
60- cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程（默认npu模式）60+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程（默认npu模式）
61 python3 ../scripts/gen_data.py # 生成测试输入数据61 python3 ../scripts/gen_data.py # 生成测试输入数据
62 ./demo # 执行编译生成的可执行程序，执行样例62 ./demo # 执行编译生成的可执行程序，执行样例
63 ```63 ```
@@ -74,10 +74,10 @@
74 74 
75- 编译选项说明75- 编译选项说明
76 76 
77-| 选项 | 可选值 | 说明 |77+| 选项　　　 | 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
78-|------|--------|------|78+| ------------| -----------------------------| ---------------------------------------------------|
79-| `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |79+| `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
80-| `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认） | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |80+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
81 81 
82- 执行结果82- 执行结果
83 83 
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/mul/data_utils.h+1-1
@@ -11,7 +11,7 @@
11 11 
12/* !12/* !
13 * \file data_utils.h13 * \file data_utils.h
14- * \brief14+ * \brief Mul样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15 */15 */
16 16 
17#ifndef DATA_UTILS_H17#ifndef DATA_UTILS_H
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/mul/mul.asc+6-0
@@ -9,6 +9,11 @@
9*/9*/
10 10 
11 11 
12+/* !
13+ * \file mul.asc
14+ * \brief 演示使用Mul API实现向量逐元素乘法计算功能，对两个输入向量逐元素相乘并输出结果
15+ */
16+ 
12#include "acl/acl.h"17#include "acl/acl.h"
13#include "kernel_operator.h"18#include "kernel_operator.h"
14#include "data_utils.h"19#include "data_utils.h"
@@ -82,6 +87,7 @@ private:
82 87 
83__global__ __vector__ void mul(GM_ADDR x, GM_ADDR y, GM_ADDR z)88__global__ __vector__ void mul(GM_ADDR x, GM_ADDR y, GM_ADDR z)
84{89{
90+ AscendC::InitSocState();
85 uint32_t totalLength = 256;91 uint32_t totalLength = 256;
86 AscendC::InitSocState();92 AscendC::InitSocState();
87 Kernel<float> kernel;93 Kernel<float> kernel;
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/muls/README.md+4-4
@@ -78,10 +78,10 @@
78 78 
79- 编译选项说明79- 编译选项说明
80 80 
81-| 选项 | 可选值 | 说明 |81+| 选项　　　　　　　　　　　| 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
82-|------|--------|------|82+| ---------------------------| -----------------------------| ---------------------------------------------------|
83-| `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |83+| `CMAKE_ASC_RUN_MODE`　　　| `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
84-| `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认） | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |84+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认）　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
85 85 
86- 执行结果86- 执行结果
87 87 
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/muls/data_utils.h+1-1
@@ -11,7 +11,7 @@
11 11 
12/* !12/* !
13 * \file data_utils.h13 * \file data_utils.h
14- * \brief14+ * \brief Muls样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15 */15 */
16 16 
17#ifndef DATA_UTILS_H17#ifndef DATA_UTILS_H
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/muls/muls.asc+6-1
@@ -1,5 +1,5 @@
1/**1/**
2-* Copyright (c) 2025 Huawei Technologies Co., Ltd.2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3* This program is free software, you can redistribute it and/or modify it under the terms and conditions of3* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4* CANN Open Software License Agreement Version 2.0 (the "License").4* CANN Open Software License Agreement Version 2.0 (the "License").
5* Please refer to the License for details. You may not use this file except in compliance with the License.5* Please refer to the License for details. You may not use this file except in compliance with the License.
@@ -9,6 +9,11 @@
9*/9*/
10 10 
11 11 
12+/* !
13+ * \file muls.asc
14+ * \brief 演示使用Muls API实现向量与标量乘法计算功能，将输入向量每个元素与标量相乘并输出结果
15+ */
16+ 
12#include "acl/acl.h"17#include "acl/acl.h"
13#include "kernel_operator.h"18#include "kernel_operator.h"
14#include "data_utils.h"19#include "data_utils.h"
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce/CMakeLists.txt+35-0
@@ -0,0 +1,35 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
cellur_z
cellur_zcellur_z4月16日

新增的文件是不是要写2026

likedislike
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+cmake_minimum_required(VERSION 3.16)
13+ 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
16+ 
17+find_package(ASC REQUIRED)
18+ 
19+project(kernel_samples LANGUAGES ASC CXX)
20+ 
21+add_executable(demo
22+ reduce.asc
23+)
24+ 
25+target_link_libraries(demo PRIVATE
26+ tiling_api
27+ register
28+ platform
29+ m
30+)
31+ 
32+target_compile_options(demo PRIVATE
33+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
34+ $<$<COMPILE_LANGUAGE:ASC>:--cce-simd-vf-fusion=false>
35+)
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce/README.md+93-0
@@ -0,0 +1,93 @@
1+# reduce样例
2+ 
3+## 概述
4+本样例基于RegBase编程范式实现Reduce运算，主要调用Reduce接口（SUM模式）。
5+- Reduce接口支持SUM/MAX/MIN归约模式，本样例以SUM模式为例
6+ 
7+## 支持的产品
8+- Ascend 950PR/Ascend 950DT
9+ 
10+## 目录结构介绍
11+```
12+├── reduce
13+│ ├── scripts
14+│ │ ├── gen_data.py // 输入数据和真值数据生成脚本
15+│ ├── CMakeLists.txt // 编译工程文件
16+│ ├── data_utils.h // 数据读入写出函数
17+│ ├── reduce.asc // AscendC样例实现 & 调用样例
18+│ └── README.md // 样例介绍
19+```
20+ 
21+## 样例描述
22+- 样例功能：
23+ 对输入向量做归约求和运算，向量shape为[1, 256]，数据类型为float，输出为归约求和结果。
24+- 样例规格：
25+ <table>
26+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="3" align="center">AIV样例</td></tr>
27+ </tr>
28+ <tr><td rowspan="2" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
29+ <tr><td align="center">x</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
30+ </tr>
31+ <tr><td rowspan="1" align="center">样例输出</td><td align="center">y</td><td align="center">[1]</td><td align="center">float</td></tr>
32+ </tr>
33+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">reduce</td></tr>
34+ </table>
35+- 样例实现：
36+ ReduceSumVF函数内调用Reduce接口进行归约计算：
37+ - 使用LoadAlign加载数据到寄存器
38+ - 使用Reduce(ReduceType::SUM)对单个repeat内所有元素求和
39+ - 使用Add累加多个repeat的局部和
40+ - 使用StoreAlign将最终结果写回UB
41+ - 调用实现
42+ 使用内核调用符<<<>>>调用核函数。
43+ 
44+## 编译运行
45+在本样例根目录下执行如下步骤，编译并执行样例。
46+- 配置环境变量
47+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
48+ - 默认路径，root用户安装CANN软件包
49+ ```bash
50+ source /usr/local/Ascend/cann/set_env.sh
51+ ```
52+ 
53+ - 默认路径，非root用户安装CANN软件包
54+ ```bash
55+ source $HOME/Ascend/cann/set_env.sh
56+ ```
57+ 
58+ - 指定路径install_path，安装CANN软件包
59+ ```bash
60+ source ${install_path}/cann/set_env.sh
61+ ```
62+ 
63+- 样例执行
64+ ```bash
65+ mkdir -p build && cd build; # 创建并进入build目录
66+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程（默认npu模式）
67+ python3 ../scripts/gen_data.py # 生成测试输入数据
68+ ./demo # 执行编译生成的可执行程序，执行样例
69+ ```
70+ 
71+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
72+ 
73+ 示例如下：
74+ ```bash
75+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
76+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
77+ ```
78+ 
79+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
80+ 
81+- 编译选项说明
82+ 
83+| 选项　　　 | 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
84+| ------------| -----------------------------| ---------------------------------------------------|
85+| `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
86+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
87+ 
88+- 执行结果
89+ 
90+ 执行结果如下，说明精度对比成功。
91+ ```bash
92+ test pass!
93+ ```
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce/data_utils.h+94-0
@@ -0,0 +1,94 @@
1+/**
2+* Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file data_utils.h
14+ * \brief ReduceSum样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15+ */
16+ 
17+#ifndef DATA_UTILS_H
18+#define DATA_UTILS_H
19+#include <fcntl.h>
20+#include <sys/stat.h>
21+#include <unistd.h>
22+#include <fstream>
23+ 
24+#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
25+ 
26+bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize)
27+{
28+ struct stat sBuf;
29+ int fileStatus = stat(filePath.data(), &sBuf);
30+ if (fileStatus == -1) {
31+ ERROR_LOG("failed to get file");
32+ return false;
33+ }
34+ if (S_ISREG(sBuf.st_mode) == 0) {
35+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
36+ return false;
37+ }
38+ 
39+ std::ifstream file;
40+ file.open(filePath, std::ios::binary);
41+ if (!file.is_open()) {
42+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
43+ return false;
44+ }
45+ 
46+ std::filebuf *buf = file.rdbuf();
47+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
48+ if (size == 0) {
49+ ERROR_LOG("file size is 0");
50+ file.close();
51+ return false;
52+ }
53+ if (size > bufferSize) {
54+ ERROR_LOG("file size is larger than buffer size");
55+ file.close();
56+ return false;
57+ }
58+ buf->pubseekpos(0, std::ios::in);
59+ buf->sgetn(static_cast<char *>(buffer), size);
60+ fileSize = size;
61+ file.close();
62+ return true;
63+}
64+ 
65+/**
66+ * @brief Write data to file
67+ * @param [in] filePath: file path
68+ * @param [in] buffer: data to write to file
69+ * @param [in] size: size to write
70+ * @return write result
71+ */
72+bool WriteFile(const std::string &filePath, const void *buffer, size_t size)
73+{
74+ if (buffer == nullptr) {
75+ ERROR_LOG("Write file failed. buffer is nullptr");
76+ return false;
77+ }
78+ 
79+ int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE);
80+ if (fd < 0) {
81+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
82+ return false;
83+ }
84+ 
85+ size_t writeSize = write(fd, buffer, size);
86+ (void)close(fd);
87+ if (writeSize != size) {
88+ ERROR_LOG("Write file Failed.");
89+ return false;
90+ }
91+ 
92+ return true;
93+}
94+#endif // DATA_UTILS_H
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce/reduce.asc+169-0
@@ -0,0 +1,169 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file reduce.asc
14+ * \brief 演示使用ReduceSum API实现向量归约求和计算功能，对输入向量所有元素求和并输出标量结果
15+ */
16+ 
17+#include "acl/acl.h"
18+#include "kernel_operator.h"
19+#include "data_utils.h"
20+ 
21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
25+template <typename T>
26+__simd_vf__ inline void ReduceSumVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr,
27+ uint32_t totalCount, uint32_t repeatTimes, uint32_t oneRepeatSize)
28+{
29+ AscendC::Reg::MaskReg mask = AscendC::Reg::CreateMask<T, AscendC::Reg::MaskPattern::ALL>();
30+ AscendC::Reg::RegTensor<T> srcReg;
31+ AscendC::Reg::RegTensor<T> dstReg;
32+ AscendC::Reg::RegTensor<T> accReg;
33+ // 初始化累加寄存器为0
34+ AscendC::Reg::Duplicate(accReg, (T)0, mask);
35+ for (uint16_t i = 0; i < repeatTimes; ++i) {
36+ AscendC::Reg::LoadAlign(srcReg, srcAddr + i * oneRepeatSize);
37+ AscendC::Reg::Reduce<AscendC::Reg::ReduceType::SUM>(dstReg, srcReg, mask);
38+ // 直接用dstReg作为Add源，无需存回UB
39+ AscendC::Reg::Add(accReg, accReg, dstReg, mask);
40+ }
41+ // 将最终累加结果写入目的地址
42+ AscendC::Reg::StoreAlign(dstAddr, accReg, mask);
43+}
44+ 
45+template <typename T>
46+class Kernel {
47+public:
48+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t totalLength)
49+ {
50+ this->totalLength = totalLength;
51+ this->xGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(x));
52+ this->yGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(y));
53+ }
54+ __aicore__ inline void Process()
55+ {
56+ AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator;
57+ AscendC::LocalTensor<T> xLocal = ubAllocator.Alloc<T, 256>();
58+ AscendC::LocalTensor<T> yLocal = ubAllocator.Alloc<T, 256>();
59+ 
60+ // CopyIn: GM -> UB
61+ AscendC::DataCopy(xLocal, xGm, totalLength);
62+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
63+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
64+ 
65+ // Compute: RegBase ReduceSum
66+ __ubuf__ T* xAddr = reinterpret_cast<__ubuf__ T*>(xLocal.GetPhyAddr());
67+ __ubuf__ T* yAddr = reinterpret_cast<__ubuf__ T*>(yLocal.GetPhyAddr());
68+ // 一次repeat计算的数据个数
69+ constexpr uint32_t oneRepeatSize = AscendC::GetVecLen() / sizeof(T);
70+ // 总迭代次数
71+ uint16_t repeatTimes = AscendC::CeilDivision(totalLength, oneRepeatSize);
72+ ReduceSumVF(yAddr, xAddr, totalLength, repeatTimes, oneRepeatSize);
73+ 
74+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
75+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
76+ 
77+ // CopyOut: UB -> GM (DataCopy要求32B对齐，float需至少8个元素)
78+ constexpr uint32_t outAlignSize = 8;
79+ AscendC::DataCopy(yGm, yLocal, outAlignSize);
80+ }
81+ 
82+private:
83+ uint32_t totalLength;
84+ AscendC::GlobalTensor<T> xGm;
85+ AscendC::GlobalTensor<T> yGm;
86+};
87+ 
88+__global__ __vector__ void reduce(GM_ADDR x, GM_ADDR y)
89+{
90+ AscendC::InitSocState();
91+ uint32_t totalLength = 256;
92+ Kernel<float> kernel;
93+ kernel.Init(x, y, totalLength);
94+ kernel.Process();
95+ AscendC::PipeBarrier<PIPE_ALL>();
96+}
97+ 
98+static bool CompareResult(const void* outputData, uint32_t outSize)
99+{
100+ void* goldenData;
101+ aclrtMallocHost((void**)(&goldenData), outSize);
102+ size_t goldenSize = outSize;
103+ bool ret = ReadFile("./output/golden.bin", goldenSize, goldenData, goldenSize);
104+ if (ret) {
105+ printf("ReadFile golden.bin success!\n");
106+ } else {
107+ printf("test failed!\n");
108+ return false;
109+ }
110+ constexpr float EPS = 1e-4;
111+ int64_t wrongNum = 0;
112+ 
113+ // Reduce结果在第一个元素，只比较第一个元素
114+ float a = (reinterpret_cast<const float*>(outputData))[0];
115+ float b = (reinterpret_cast<const float*>(goldenData))[0];
116+ float ae = std::abs(a - b);
117+ float re = ae / std::abs(b);
118+ if (ae > EPS && re > EPS) {
119+ printf("CompareResult golden.bin failed output is %lf, golden is %lf\n", a, b);
120+ wrongNum++;
121+ }
122+ aclrtFreeHost(goldenData);
123+ if (wrongNum != 0) {
124+ return false;
125+ } else {
126+ printf("CompareResult golden.bin success!\n");
127+ return true;
128+ }
129+}
130+ 
131+int32_t main(int32_t argc, char* argv[])
132+{
133+ uint32_t numBlocks = 1;
134+ aclInit(nullptr);
135+ int32_t deviceId = 0;
136+ aclrtSetDevice(deviceId);
137+ aclrtStream stream = nullptr;
138+ aclrtCreateStream(&stream);
139+ uint32_t totalLength = 256;
140+ size_t inputByteSize = totalLength * sizeof(float);
141+ size_t outputByteSize = 8 * sizeof(float); // 32B对齐，至少8个float
142+ uint8_t *xHost, *zHost;
143+ uint8_t *xDevice, *zDevice;
144+ aclrtMallocHost((void**)(&xHost), inputByteSize);
145+ aclrtMallocHost((void**)(&zHost), outputByteSize);
146+ aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
147+ aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
148+ ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);
149+ aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
150+ reduce<<<numBlocks, nullptr, stream>>>(xDevice, zDevice);
151+ aclrtSynchronizeStream(stream);
152+ aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
153+ WriteFile("./output/output.bin", zHost, outputByteSize);
154+ bool goldenResult = true;
155+ goldenResult = CompareResult(zHost, outputByteSize);
156+ if (goldenResult) {
157+ printf("test pass!\n");
158+ } else {
159+ printf("test failed!\n");
160+ }
161+ aclrtFree(xDevice);
162+ aclrtFree(zDevice);
163+ aclrtFreeHost(xHost);
164+ aclrtFreeHost(zHost);
165+ aclrtDestroyStream(stream);
166+ aclrtResetDevice(deviceId);
167+ aclFinalize();
168+ return 0;
169+}
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce/scripts/gen_data.py+32-0
@@ -0,0 +1,32 @@
1+#!/usr/bin/python3
2+# coding=utf-8
3+ 
4+# ----------------------------------------------------------------------------------------------------------
5+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
7+# CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
10+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11+# See LICENSE in the root of the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------------------------------------
13+ 
14+ 
15+import os
16+import numpy as np
17+ 
18+def gen_golden_data_simple():
19+ total_length = 256
20+ data_type = np.float32
21+ # 生成一个 [1, 256] 向量
22+ x = np.random.uniform(1, 2, [1, total_length]).astype(data_type)
23+ # 计算归约求和 golden = sum(x)，补齐到8个float以满足32B对齐
24+ golden = np.zeros(8, dtype=data_type)
25+ golden[0] = np.sum(x)
26+ os.makedirs("input", exist_ok=True)
27+ os.makedirs("output", exist_ok=True)
28+ x.tofile('./input/input_x.bin')
29+ golden.tofile('./output/golden.bin')
30+ 
31+if __name__ == "__main__":
32+ gen_golden_data_simple()
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_block/CMakeLists.txt+35-0
@@ -0,0 +1,35 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You may not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+cmake_minimum_required(VERSION 3.16)
13+ 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
16+ 
17+find_package(ASC REQUIRED)
18+ 
19+project(kernel_samples LANGUAGES ASC CXX)
20+ 
21+add_executable(demo
22+ reduce_block.asc
23+)
24+ 
25+target_link_libraries(demo PRIVATE
26+ tiling_api
27+ register
28+ platform
29+ m
30+)
31+ 
32+target_compile_options(demo PRIVATE
33+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
34+ $<$<COMPILE_LANGUAGE:ASC>:--cce-simd-vf-fusion=false>
35+)
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_block/README.md+98-0
@@ -0,0 +1,98 @@
1+# reduce_block样例
2+ 
3+## 概述
4+本样例基于RegBase编程范式实现ReduceDataBlock运算，主要调用ReduceDataBlock接口（SUM模式）。
5+- ReduceDataBlock接口支持SUM/MAX/MIN归约模式，本样例以SUM模式为例
6+ 
7+## 支持的产品
8+- Ascend 950PR/Ascend 950DT
9+ 
10+## 目录结构介绍
11+```
12+├── reduce_block
13+│ ├── scripts
14+│ │ ├── gen_data.py // 输入数据和真值数据生成脚本
15+│ ├── CMakeLists.txt // 编译工程文件
16+│ ├── data_utils.h // 数据读入写出函数
17+│ ├── reduce_block.asc // AscendC样例实现 & 调用样例
18+│ └── README.md // 样例介绍
19+```
20+ 
21+## 样例描述
22+- 样例功能：
23+ 对输入向量做ReduceDataBlock归约求和运算，向量shape为[1, 256]，数据类型为float。
24+ 每8个float（一个DataBlock=32B）产生一个求和结果，共输出32个结果。
25+- 每个 `oneRepeatSize`（如 float 类型为 64 元素）块中：
26+ - 前 `blocksPerRepeat`（8个）位置存放有效的归约结果
27+ - 后 `oneRepeatSize - blocksPerRepeat`（56个）位置为无效数据（未定义值）
28+ 
29+- 样例规格：
30+ <table>
31+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="3" align="center">AIV样例</td></tr>
32+ </tr>
33+ <tr><td rowspan="2" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
34+ <tr><td align="center">x</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
35+ </tr>
36+ <tr><td rowspan="1" align="center">样例输出</td><td align="center">y</td><td align="center">[1, 32]</td><td align="center">float</td></tr>
37+ </tr>
38+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">reduce_block</td></tr>
39+ </table>
40+- 样例实现：
41+ ReduceDataBlockSumVF函数内调用ReduceDataBlock接口进行归约计算：
42+ - 使用LoadAlign加载数据到寄存器
43+ - 使用ReduceDataBlock(ReduceType::SUM)对每个DataBlock(32B)内的元素求和
44+ - 结果采用交织布局存储在dstReg中，每个repeat的前8个位置为有效结果
45+ - 使用StoreAlign将结果写回UB，偏移使用oneRepeatSize
46+ - 调用实现
47+ 使用内核调用符<<<>>>调用核函数。
48+ 
49+## 编译运行
50+在本样例根目录下执行如下步骤，编译并执行样例。
51+- 配置环境变量
52+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
53+ - 默认路径，root用户安装CANN软件包
54+ ```bash
55+ source /usr/local/Ascend/cann/set_env.sh
56+ ```
57+ 
58+ - 默认路径，非root用户安装CANN软件包
59+ ```bash
60+ source $HOME/Ascend/cann/set_env.sh
61+ ```
62+ 
63+ - 指定路径install_path，安装CANN软件包
64+ ```bash
65+ source ${install_path}/cann/set_env.sh
66+ ```
67+ 
68+- 样例执行
69+ ```bash
70+ mkdir -p build && cd build; # 创建并进入build目录
71+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程（默认npu模式）
72+ python3 ../scripts/gen_data.py # 生成测试输入数据
73+ ./demo # 执行编译生成的可执行程序，执行样例
74+ ```
75+ 
76+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
77+ 
78+ 示例如下：
79+ ```bash
80+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
81+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
82+ ```
83+ 
84+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
85+ 
86+- 编译选项说明
87+ 
88+| 选项　　　 | 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
89+| ------------| -----------------------------| ---------------------------------------------------|
90+| `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
91+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
92+ 
93+- 执行结果
94+ 
95+ 执行结果如下，说明精度对比成功。
96+ ```bash
97+ test pass!
98+ ```
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_block/data_utils.h+94-0
@@ -0,0 +1,94 @@
1+/**
2+* Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file data_utils.h
14+ * \brief BlockReduceSum样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15+ */
16+ 
17+#ifndef DATA_UTILS_H
18+#define DATA_UTILS_H
19+#include <fcntl.h>
20+#include <sys/stat.h>
21+#include <unistd.h>
22+#include <fstream>
23+ 
24+#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
25+ 
26+bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize)
27+{
28+ struct stat sBuf;
29+ int fileStatus = stat(filePath.data(), &sBuf);
30+ if (fileStatus == -1) {
31+ ERROR_LOG("failed to get file");
32+ return false;
33+ }
34+ if (S_ISREG(sBuf.st_mode) == 0) {
35+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
36+ return false;
37+ }
38+ 
39+ std::ifstream file;
40+ file.open(filePath, std::ios::binary);
41+ if (!file.is_open()) {
42+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
43+ return false;
44+ }
45+ 
46+ std::filebuf *buf = file.rdbuf();
47+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
48+ if (size == 0) {
49+ ERROR_LOG("file size is 0");
50+ file.close();
51+ return false;
52+ }
53+ if (size > bufferSize) {
54+ ERROR_LOG("file size is larger than buffer size");
55+ file.close();
56+ return false;
57+ }
58+ buf->pubseekpos(0, std::ios::in);
59+ buf->sgetn(static_cast<char *>(buffer), size);
60+ fileSize = size;
61+ file.close();
62+ return true;
63+}
64+ 
65+/**
66+ * @brief Write data to file
67+ * @param [in] filePath: file path
68+ * @param [in] buffer: data to write to file
69+ * @param [in] size: size to write
70+ * @return write result
71+ */
72+bool WriteFile(const std::string &filePath, const void *buffer, size_t size)
73+{
74+ if (buffer == nullptr) {
75+ ERROR_LOG("Write file failed. buffer is nullptr");
76+ return false;
77+ }
78+ 
79+ int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE);
80+ if (fd < 0) {
81+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
82+ return false;
83+ }
84+ 
85+ size_t writeSize = write(fd, buffer, size);
86+ (void)close(fd);
87+ if (writeSize != size) {
88+ ERROR_LOG("Write file Failed.");
89+ return false;
90+ }
91+ 
92+ return true;
93+}
94+#endif // DATA_UTILS_H
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_block/reduce_block.asc+192-0
@@ -0,0 +1,192 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You may not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file reduce_block.asc
14+ * \brief 演示使用BlockReduceSum API实现块内归约求和计算功能，对每个数据块内的元素求和并输出结果
15+ */
16+ 
17+#include "acl/acl.h"
18+#include "kernel_operator.h"
19+#include "data_utils.h"
20+ 
21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
25+template <typename T>
26+__simd_vf__ inline void ReduceDataBlockSumVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr,
27+ uint32_t repeatTimes, uint32_t oneRepeatSize)
28+{
29+ AscendC::Reg::RegTensor<T> srcReg;
30+ AscendC::Reg::RegTensor<T> dstReg;
31+ AscendC::Reg::MaskReg mask = AscendC::Reg::CreateMask<T, AscendC::Reg::MaskPattern::ALL>();
32+ 
33+ for (uint16_t i = 0; i < repeatTimes; ++i) {
34+ uint32_t offset = i * oneRepeatSize;
35+ AscendC::Reg::LoadAlign(srcReg, srcAddr + offset);
36+ // ReduceDataBlock 输出依次存储在dstReg的最低位
37+ AscendC::Reg::ReduceDataBlock<AscendC::Reg::ReduceType::SUM>(dstReg, srcReg, mask);
38+ // 使用 StoreAlign，偏移使用 oneRepeatSize
39+ AscendC::Reg::StoreAlign(dstAddr + i * oneRepeatSize, dstReg, mask);
40+ }
41+}
42+ 
43+template <typename T>
44+class Kernel {
45+public:
46+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t totalLength)
47+ {
48+ this->totalLength = totalLength;
49+ this->xGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(x));
50+ this->yGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(y));
51+ }
52+ __aicore__ inline void Process()
53+ {
54+ AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator;
55+ AscendC::LocalTensor<T> xLocal = ubAllocator.Alloc<T, 256>();
56+ // 输出需要 repeatTimes * oneRepeatSize = 4 * 64 = 256 个 float
57+ // ReduceDataBlock 采用交织布局，每次repeat存储oneRepeatSize个元素
58+ // 其中前8个有效，后56个为无效数据
59+ AscendC::LocalTensor<T> yLocal = ubAllocator.Alloc<T, 256>();
60+ 
61+ // CopyIn: GM -> UB
62+ AscendC::DataCopy(xLocal, xGm, totalLength);
63+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
64+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
65+ 
66+ // Compute: RegBase ReduceDataBlock SUM
67+ __ubuf__ T* xAddr = reinterpret_cast<__ubuf__ T*>(xLocal.GetPhyAddr());
68+ __ubuf__ T* yAddr = reinterpret_cast<__ubuf__ T*>(yLocal.GetPhyAddr());
69+ // 一次 repeat 计算的数据个数
70+ constexpr uint32_t oneRepeatSize = AscendC::GetVecLen() / sizeof(T);
71+ // 总迭代次数
72+ uint16_t repeatTimes = AscendC::CeilDivision(totalLength, oneRepeatSize);
73+ ReduceDataBlockSumVF(yAddr, xAddr, repeatTimes, oneRepeatSize);
74+ 
75+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
76+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
77+ 
78+ // CopyOut: UB -> GM
79+ // ReduceDataBlock输出: 使用交织布局
80+ // 每次repeat存储oneRepeatSize个元素，其中前8个有效
81+ // 总输出大小 = repeatTimes * oneRepeatSize
82+ uint32_t outputSize = repeatTimes * oneRepeatSize;
83+ AscendC::DataCopy(yGm, yLocal, outputSize);
84+ }
85+ 
86+private:
87+ uint32_t totalLength;
88+ AscendC::GlobalTensor<T> xGm;
89+ AscendC::GlobalTensor<T> yGm;
90+};
91+ 
92+__global__ __vector__ void reduce_block(GM_ADDR x, GM_ADDR y)
93+{
94+ AscendC::InitSocState();
95+ uint32_t totalLength = 256;
96+ Kernel<float> kernel;
97+ kernel.Init(x, y, totalLength);
98+ kernel.Process();
99+ AscendC::PipeBarrier<PIPE_ALL>();
100+}
101+ 
102+static bool CompareResult(const void* outputData, uint32_t outSize, uint32_t totalLength)
103+{
104+ void* goldenData;
105+ aclrtMallocHost((void**)(&goldenData), outSize);
106+ size_t goldenSize = outSize;
107+ bool ret = ReadFile("./output/golden.bin", goldenSize, goldenData, goldenSize);
108+ if (ret) {
109+ printf("ReadFile golden.bin success!\n");
110+ } else {
111+ printf("test failed!\n");
112+ return false;
113+ }
114+ constexpr float EPS = 1e-4;
115+ int64_t wrongNum = 0;
116+ 
117+ // ReduceDataBlock: 结果采用交织布局
118+ // 每次repeat存储oneRepeatSize个元素，其中前blocksPerRepeat(8)个有效
119+ // 读取第i个有效结果时，需按交织布局计算实际索引
120+ constexpr uint32_t oneRepeatSize = 64; // GetVecLen() / sizeof(float) = 64
121+ constexpr uint32_t blocksPerRepeat = 8; // oneRepeatSize / 8 = 8 (每个DataBlock=8个float)
122+ uint32_t numResults = totalLength / 8;
123+ 
124+ const float* output = reinterpret_cast<const float*>(outputData);
125+ const float* golden = reinterpret_cast<const float*>(goldenData);
126+ 
127+ for (uint32_t i = 0; i < numResults; i++) {
128+ // 计算交织布局中的实际位置
129+ uint32_t repeatIdx = i / blocksPerRepeat;
130+ uint32_t elemIdx = i % blocksPerRepeat;
131+ uint32_t actualIdx = repeatIdx * oneRepeatSize + elemIdx;
132+ float a = output[actualIdx];
133+ float b = golden[i];
134+ float ae = std::abs(a - b);
135+ float re = ae / std::abs(b);
136+ if (ae > EPS && re > EPS) {
137+ printf("CompareResult golden.bin failed at block %zu: output is %lf, golden is %lf\n", (size_t)i, a, b);
138+ wrongNum++;
139+ }
140+ }
141+ aclrtFreeHost(goldenData);
142+ if (wrongNum != 0) {
143+ return false;
144+ } else {
145+ printf("CompareResult golden.bin success!\n");
146+ return true;
147+ }
148+}
149+ 
150+int32_t main(int32_t argc, char* argv[])
151+{
152+ uint32_t numBlocks = 1;
153+ aclInit(nullptr);
154+ int32_t deviceId = 0;
155+ aclrtSetDevice(deviceId);
156+ aclrtStream stream = nullptr;
157+ aclrtCreateStream(&stream);
158+ uint32_t totalLength = 256;
159+ size_t inputByteSize = totalLength * sizeof(float);
160+ // 输出: ReduceDataBlock使用交织布局
161+ // 每次repeat存储oneRepeatSize个元素（其中前blocksPerRepeat个有效）
162+ constexpr uint32_t oneRepeatSize = 64;
163+ uint32_t repeatTimes = (totalLength + oneRepeatSize - 1) / oneRepeatSize;
164+ size_t outputByteSize = repeatTimes * oneRepeatSize * sizeof(float);
165+ uint8_t *xHost, *zHost;
166+ uint8_t *xDevice, *zDevice;
167+ aclrtMallocHost((void**)(&xHost), inputByteSize);
168+ aclrtMallocHost((void**)(&zHost), outputByteSize);
169+ aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
170+ aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
171+ ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);
172+ aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
173+ reduce_block<<<numBlocks, nullptr, stream>>>(xDevice, zDevice);
174+ aclrtSynchronizeStream(stream);
175+ aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
176+ WriteFile("./output/output.bin", zHost, outputByteSize);
177+ bool goldenResult = true;
178+ goldenResult = CompareResult(zHost, outputByteSize, totalLength);
179+ if (goldenResult) {
180+ printf("test pass!\n");
181+ } else {
182+ printf("test failed!\n");
183+ }
184+ aclrtFree(xDevice);
185+ aclrtFree(zDevice);
186+ aclrtFreeHost(xHost);
187+ aclrtFreeHost(zHost);
188+ aclrtDestroyStream(stream);
189+ aclrtResetDevice(deviceId);
190+ aclFinalize();
191+ return 0;
192+}
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_block/scripts/gen_data.py+35-0
@@ -0,0 +1,35 @@
1+#!/usr/bin/python3
2+# coding=utf-8
3+ 
4+# ----------------------------------------------------------------------------------------------------------
5+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
7+# CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You may not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
10+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11+# See LICENSE in the root of the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------------------------------------
13+ 
14+ 
15+import os
16+import numpy as np
17+ 
18+def gen_golden_data_simple():
19+ total_length = 256
20+ data_type = np.float32
21+ # 生成一个 [1, 256] 向量
22+ x = np.random.uniform(1, 2, [1, total_length]).astype(data_type)
23+ # ReduceDataBlock: 每8个float(DataBlock=32B)求和产生一个结果
24+ # 结果依次存储在dstReg的最低位，即位置0,1,2,...,31
25+ num_blocks = total_length // 8
26+ golden = np.zeros(num_blocks, dtype=data_type)
27+ for i in range(num_blocks):
28+ golden[i] = np.sum(x[0, i*8:(i+1)*8])
29+ os.makedirs("input", exist_ok=True)
30+ os.makedirs("output", exist_ok=True)
31+ x.tofile('./input/input_x.bin')
32+ golden.tofile('./output/golden.bin')
33+ 
34+if __name__ == "__main__":
35+ gen_golden_data_simple()
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_pair/CMakeLists.txt+35-0
@@ -0,0 +1,35 @@
1+# ----------------------------------------------------------------------------------------------------------
2+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+# CANN Open Software License Agreement Version 2.0 (the "License").
5+# Please refer to the License for details. You can not use this file except in compliance with the License.
6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+# See LICENSE in the root of the software repository for the full text of the License.
9+# ----------------------------------------------------------------------------------------------------------
10+ 
11+ 
12+cmake_minimum_required(VERSION 3.16)
13+ 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
16+ 
17+find_package(ASC REQUIRED)
18+ 
19+project(kernel_samples LANGUAGES ASC CXX)
20+ 
21+add_executable(demo
22+ reduce_pair.asc
23+)
24+ 
25+target_link_libraries(demo PRIVATE
26+ tiling_api
27+ register
28+ platform
29+ m
30+)
31+ 
32+target_compile_options(demo PRIVATE
33+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
34+ $<$<COMPILE_LANGUAGE:ASC>:--cce-simd-vf-fusion=false>
35+)
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_pair/README.md+96-0
@@ -0,0 +1,96 @@
1+# reduce_pair样例
2+ 
3+## 概述
4+本样例基于RegBase编程范式实现PairReduceElem运算，主要调用PairReduceElem接口（SUM模式）。
5+- PairReduceElem接口对相邻的奇偶元素对进行归约求和，当前仅支持SUM归约模式
6+ 
7+ 
8+## 支持的产品
9+- Ascend 950PR/Ascend 950DT
10+ 
11+## 目录结构介绍
12+```
13+├── reduce_pair
14+│ ├── scripts
15+│ │ ├── gen_data.py // 输入数据和真值数据生成脚本
16+│ ├── CMakeLists.txt // 编译工程文件
17+│ ├── data_utils.h // 数据读入写出函数
18+│ ├── reduce_pair.asc // AscendC样例实现 & 调用样例
19+│ └── README.md // 样例介绍
20+```
21+ 
22+## 样例描述
23+- 样例功能：
24+ 对输入向量做PairReduceElem归约求和运算，向量shape为[1, 256]，数据类型为float。
25+ PairReduceElem对相邻奇偶元素对求和：dst[i] = src[2i] + src[2i+1]
26+ 256个float输入产生128个求和结果。
27+ 
28+- 样例规格：
29+ <table>
30+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="3" align="center">AIV样例</td></tr>
31+ </tr>
32+ <tr><td rowspan="2" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td></tr>
33+ <tr><td align="center">x</td><td align="center">[1, 256]</td><td align="center">float</td></tr>
34+ </tr>
35+ <tr><td rowspan="1" align="center">样例输出</td><td align="center">y</td><td align="center">[1, 128]</td><td align="center">float</td></tr>
36+ </tr>
37+ <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">reduce_pair</td></tr>
38+ </table>
39+- 样例实现：
40+ PairReduceSumVF函数内调用PairReduceElem接口进行归约计算：
41+ - 使用LoadAlign加载数据到寄存器
42+ - 使用PairReduceElem(PairReduce::SUM)对相邻奇偶元素对求和
43+ - 使用StoreAlign将结果写回UB
44+ - 调用实现
45+ 使用内核调用符<<<>>>调用核函数。
46+ 
47+## 编译运行
48+在本样例根目录下执行如下步骤，编译并执行样例。
49+- 配置环境变量
50+ 请根据当前环境上CANN开发套件包的[安装方式](../../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
51+ - 默认路径，root用户安装CANN软件包
52+ ```bash
53+ source /usr/local/Ascend/cann/set_env.sh
54+ ```
55+ 
56+ - 默认路径，非root用户安装CANN软件包
57+ ```bash
58+ source $HOME/Ascend/cann/set_env.sh
59+ ```
60+ 
61+ - 指定路径install_path，安装CANN软件包
62+ ```bash
63+ source ${install_path}/cann/set_env.sh
64+ ```
65+ 
66+- 样例执行
67+ ```bash
68+ mkdir -p build && cd build; # 创建并进入build目录
69+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程（默认npu模式）
70+ python3 ../scripts/gen_data.py # 生成测试输入数据
71+ ./demo # 执行编译生成的可执行程序，执行样例
72+ ```
73+ 
74+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
75+ 
76+ 示例如下：
77+ ```bash
78+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
79+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
80+ ```
81+ 
82+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
83+ 
84+- 编译选项说明
85+ 
86+| 选项　　　　　　　　　　　| 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
87+| ---------------------------| -----------------------------| ---------------------------------------------------|
88+| `CMAKE_ASC_RUN_MODE`　　　| `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
89+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510`　　　　　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
90+ 
91+- 执行结果
92+ 
93+ 执行结果如下，说明精度对比成功。
94+ ```bash
95+ test pass!
96+ ```
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_pair/data_utils.h+94-0
@@ -0,0 +1,94 @@
1+/**
2+* Copyright (c) 2025 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You can not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file data_utils.h
14+ * \brief PairReduceSum样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15+ */
16+ 
17+#ifndef DATA_UTILS_H
18+#define DATA_UTILS_H
19+#include <fcntl.h>
20+#include <sys/stat.h>
21+#include <unistd.h>
22+#include <fstream>
23+ 
24+#define ERROR_LOG(fmt, args...) fprintf(stdout, "[ERROR] " fmt "\n", ##args)
25+ 
26+bool ReadFile(const std::string &filePath, size_t &fileSize, void *buffer, size_t bufferSize)
27+{
28+ struct stat sBuf;
29+ int fileStatus = stat(filePath.data(), &sBuf);
30+ if (fileStatus == -1) {
31+ ERROR_LOG("failed to get file");
32+ return false;
33+ }
34+ if (S_ISREG(sBuf.st_mode) == 0) {
35+ ERROR_LOG("%s is not a file, please enter a file", filePath.c_str());
36+ return false;
37+ }
38+ 
39+ std::ifstream file;
40+ file.open(filePath, std::ios::binary);
41+ if (!file.is_open()) {
42+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
43+ return false;
44+ }
45+ 
46+ std::filebuf *buf = file.rdbuf();
47+ size_t size = buf->pubseekoff(0, std::ios::end, std::ios::in);
48+ if (size == 0) {
49+ ERROR_LOG("file size is 0");
50+ file.close();
51+ return false;
52+ }
53+ if (size > bufferSize) {
54+ ERROR_LOG("file size is larger than buffer size");
55+ file.close();
56+ return false;
57+ }
58+ buf->pubseekpos(0, std::ios::in);
59+ buf->sgetn(static_cast<char *>(buffer), size);
60+ fileSize = size;
61+ file.close();
62+ return true;
63+}
64+ 
65+/**
66+ * @brief Write data to file
67+ * @param [in] filePath: file path
68+ * @param [in] buffer: data to write to file
69+ * @param [in] size: size to write
70+ * @return write result
71+ */
72+bool WriteFile(const std::string &filePath, const void *buffer, size_t size)
73+{
74+ if (buffer == nullptr) {
75+ ERROR_LOG("Write file failed. buffer is nullptr");
76+ return false;
77+ }
78+ 
79+ int fd = open(filePath.c_str(), O_RDWR | O_CREAT | O_TRUNC, S_IRUSR | S_IWRITE);
80+ if (fd < 0) {
81+ ERROR_LOG("Open file failed. path = %s", filePath.c_str());
82+ return false;
83+ }
84+ 
85+ size_t writeSize = write(fd, buffer, size);
86+ (void)close(fd);
87+ if (writeSize != size) {
88+ ERROR_LOG("Write file Failed.");
89+ return false;
90+ }
91+ 
92+ return true;
93+}
94+#endif // DATA_UTILS_H
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_pair/reduce_pair.asc+195-0
@@ -0,0 +1,195 @@
1+/**
2+* Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+* This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+* CANN Open Software License Agreement Version 2.0 (the "License").
5+* Please refer to the License for details. You can not use this file except in compliance with the License.
6+* THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+* INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+* See LICENSE in the root of the software repository for the full text of the License.
9+*/
10+ 
11+ 
12+/* !
13+ * \file reduce_pair.asc
14+ * \brief 演示使用PairReduceSum API实现成对归约求和计算功能，对相邻元素成对求和并输出结果
15+ */
16+ 
17+#include "acl/acl.h"
18+#include "kernel_operator.h"
19+#include "data_utils.h"
20+ 
21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
25+template <typename T>
26+__simd_vf__ inline void PairReduceSumVF(__ubuf__ T* dstAddr, __ubuf__ T* srcAddr,
27+ uint32_t count, uint32_t repeatTimes,
28+ uint32_t oneRepeatSize)
29+{
30+ AscendC::Reg::RegTensor<T> srcReg;
31+ AscendC::Reg::RegTensor<T> dstReg;
32+ AscendC::Reg::MaskReg mask;
33+ 
34+ for (uint16_t i = 0; i < repeatTimes; ++i) {
35+ AscendC::Reg::LoadAlign(srcReg, srcAddr + i * oneRepeatSize);
36+ mask = AscendC::Reg::UpdateMask<T>(count);
37+ // PairReduceElem: 相邻奇偶元素对求和
38+ // dst[0] = src[0] + src[1], dst[1] = src[2] + src[3], ...
39+ AscendC::Reg::PairReduceElem<AscendC::Reg::PairReduce::SUM>(dstReg, srcReg, mask);
40+ // 使用 StoreAlign，偏移使用 oneRepeatSize
41+ // PairReduceElem 输出在寄存器的前半部分
42+ AscendC::Reg::StoreAlign(dstAddr + i * oneRepeatSize, dstReg, mask);
43+ }
44+}
45+ 
46+template <typename T>
47+class Kernel {
48+public:
49+ __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t totalLength)
50+ {
51+ this->totalLength = totalLength;
52+ this->xGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(x));
53+ this->yGm.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(y));
54+ }
55+ __aicore__ inline void Process()
56+ {
57+ AscendC::LocalMemAllocator<AscendC::Hardware::UB> ubAllocator;
58+ // 输入 256 个 float
59+ AscendC::LocalTensor<T> xLocal = ubAllocator.Alloc<T, 256>();
60+ // 输出需要 repeatTimes * oneRepeatSize = 4 * 64 = 256 个 float
61+ // 但 PairReduceElem 只产生有效数据在寄存器前半部分
62+ AscendC::LocalTensor<T> yLocal = ubAllocator.Alloc<T, 512>();
63+ 
64+ // CopyIn: GM -> UB
65+ AscendC::DataCopy(xLocal, xGm, totalLength);
66+ AscendC::SetFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
67+ AscendC::WaitFlag<AscendC::HardEvent::MTE2_V>(EVENT_ID0);
68+ 
69+ // Compute: RegBase PairReduceElem SUM
70+ __ubuf__ T* xAddr = reinterpret_cast<__ubuf__ T*>(xLocal.GetPhyAddr());
71+ __ubuf__ T* yAddr = reinterpret_cast<__ubuf__ T*>(yLocal.GetPhyAddr());
72+ // 一次repeat计算的数据个数
73+ constexpr uint32_t oneRepeatSize = AscendC::GetVecLen() / sizeof(T);
74+ // 总迭代次数
75+ uint16_t repeatTimes = AscendC::CeilDivision(totalLength, oneRepeatSize);
76+ PairReduceSumVF(yAddr, xAddr, totalLength, repeatTimes, oneRepeatSize);
77+ 
78+ AscendC::SetFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
79+ AscendC::WaitFlag<AscendC::HardEvent::V_MTE3>(EVENT_ID0);
80+ 
81+ // CopyOut: UB -> GM
82+ // PairReduceElem输出: 使用交错布局
83+ // 每次repeat存储oneRepeatSize个元素，其中前halfRepeatSize个有效
84+ // 总输出大小 = repeatTimes * oneRepeatSize
85+ uint32_t outputSize = repeatTimes * oneRepeatSize;
86+ AscendC::DataCopy(yGm, yLocal, outputSize);
87+ }
88+ 
89+private:
90+ uint32_t totalLength;
91+ AscendC::GlobalTensor<T> xGm;
92+ AscendC::GlobalTensor<T> yGm;
93+};
94+ 
95+__global__ __vector__ void reduce_pair(GM_ADDR x, GM_ADDR y)
96+{
97+ AscendC::InitSocState();
98+ uint32_t totalLength = 256;
99+ Kernel<float> kernel;
100+ kernel.Init(x, y, totalLength);
101+ kernel.Process();
102+ AscendC::PipeBarrier<PIPE_ALL>();
103+}
104+ 
105+static bool CompareResult(const void* outputData, uint32_t outSize, uint32_t totalLength)
106+{
107+ void* goldenData;
108+ aclrtMallocHost((void**)(&goldenData), outSize);
109+ size_t goldenSize = outSize;
110+ bool ret = ReadFile("./output/golden.bin", goldenSize, goldenData, goldenSize);
111+ if (ret) {
112+ printf("ReadFile golden.bin success!\n");
113+ } else {
114+ printf("test failed!\n");
115+ return false;
116+ }
117+ constexpr float EPS = 1e-4;
118+ int64_t wrongNum = 0;
119+ 
120+ // PairReduceElem: 结果数量为输入的一半
121+ // 使用交错布局：每64个元素中只有前32个有效
122+ constexpr uint32_t oneRepeatSize = 64; // GetVecLen() / sizeof(float) = 64
123+ constexpr uint32_t halfRepeatSize = 32; // oneRepeatSize / 2
124+ uint32_t numResults = totalLength / 2;
125+ uint32_t numRepeats = (numResults + halfRepeatSize - 1) / halfRepeatSize;
126+ 
127+ const float* output = reinterpret_cast<const float*>(outputData);
128+ const float* golden = reinterpret_cast<const float*>(goldenData);
129+ 
130+ for (uint32_t i = 0; i < numResults; i++) {
131+ // 计算交错布局中的实际位置
132+ uint32_t repeatIdx = i / halfRepeatSize;
133+ uint32_t elemIdx = i % halfRepeatSize;
134+ uint32_t actualIdx = repeatIdx * oneRepeatSize + elemIdx;
135+ float a = output[actualIdx];
136+ float b = golden[i];
137+ float ae = std::abs(a - b);
138+ float re = ae / std::abs(b);
139+ if (ae > EPS && re > EPS) {
140+ printf("CompareResult golden.bin failed at index %zu: output is %lf, golden is %lf\n", i, a, b);
141+ wrongNum++;
142+ }
143+ }
144+ aclrtFreeHost(goldenData);
145+ if (wrongNum != 0) {
146+ return false;
147+ } else {
148+ printf("CompareResult golden.bin success!\n");
149+ return true;
150+ }
151+}
152+ 
153+int32_t main(int32_t argc, char* argv[])
154+{
155+ uint32_t numBlocks = 1;
156+ aclInit(nullptr);
157+ int32_t deviceId = 0;
158+ aclrtSetDevice(deviceId);
159+ aclrtStream stream = nullptr;
160+ aclrtCreateStream(&stream);
161+ uint32_t totalLength = 256;
162+ size_t inputByteSize = totalLength * sizeof(float);
163+ // 输出: PairReduceElem使用交错布局
164+ // 每次repeat存储oneRepeatSize个元素（其中前halfRepeatSize个有效）
165+ constexpr uint32_t oneRepeatSize = 64;
166+ uint32_t repeatTimes = (totalLength + oneRepeatSize - 1) / oneRepeatSize;
167+ size_t outputByteSize = repeatTimes * oneRepeatSize * sizeof(float);
168+ uint8_t *xHost, *zHost;
169+ uint8_t *xDevice, *zDevice;
170+ aclrtMallocHost((void**)(&xHost), inputByteSize);
171+ aclrtMallocHost((void**)(&zHost), outputByteSize);
172+ aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
173+ aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST);
174+ ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);
175+ aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE);
176+ reduce_pair<<<numBlocks, nullptr, stream>>>(xDevice, zDevice);
177+ aclrtSynchronizeStream(stream);
178+ aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST);
179+ WriteFile("./output/output.bin", zHost, outputByteSize);
180+ bool goldenResult = true;
181+ goldenResult = CompareResult(zHost, outputByteSize, totalLength);
182+ if (goldenResult) {
183+ printf("test pass!\n");
184+ } else {
185+ printf("test failed!\n");
186+ }
187+ aclrtFree(xDevice);
188+ aclrtFree(zDevice);
189+ aclrtFreeHost(xHost);
190+ aclrtFreeHost(zHost);
191+ aclrtDestroyStream(stream);
192+ aclrtResetDevice(deviceId);
193+ aclFinalize();
194+ return 0;
195+}
Aexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/reduce_pair/scripts/gen_data.py+40-0
@@ -0,0 +1,40 @@
1+#!/usr/bin/python3
2+# coding=utf-8
3+ 
4+# ----------------------------------------------------------------------------------------------------------
5+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
6+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
7+# CANN Open Software License Agreement Version 2.0 (the "License").
8+# Please refer to the License for details. You can not use this file except in compliance with the License.
9+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
10+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
11+# See LICENSE in the root of the software repository for the full text of the License.
12+# ----------------------------------------------------------------------------------------------------------
13+ 
14+ 
15+import os
16+import numpy as np
17+ 
18+def gen_golden_data_simple():
19+ total_length = 256
20+ data_type = np.float32
21+ # 生成一个 [1, 256] 向量
22+ x = np.random.uniform(1, 2, [1, total_length]).astype(data_type)
23+ # PairReduceElem: 相邻奇偶元素对求和
24+ # golden[0] = x[0] + x[1], golden[1] = x[2] + x[3], ...
25+ # 结果数量为输入的一半
26+ golden = np.zeros(total_length // 2, dtype=data_type)
27+ for i in range(total_length // 2):
28+ golden[i] = x[0, 2 * i] + x[0, 2 * i + 1]
29+ # 补齐到8个float以满足32B对齐
30+ num_results = total_length // 2
31+ aligned_size = ((num_results + 7) // 8) * 8
32+ golden_aligned = np.zeros(aligned_size, dtype=data_type)
33+ golden_aligned[:num_results] = golden
34+ os.makedirs("input", exist_ok=True)
35+ os.makedirs("output", exist_ok=True)
36+ x.tofile('./input/input_x.bin')
37+ golden_aligned.tofile('./output/golden.bin')
38+ 
39+if __name__ == "__main__":
40+ gen_golden_data_simple()
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/truncate/data_utils.h+1-1
@@ -11,7 +11,7 @@
11 11 
12/* !12/* !
13 * \file data_utils.h13 * \file data_utils.h
14- * \brief14+ * \brief Trunc样例的输入输出数据读写辅助工具，提供文件读取、写入功能
15 */15 */
16 16 
17#ifndef DATA_UTILS_H17#ifndef DATA_UTILS_H
Mexamples/01_simd_cpp_api/02_features/03_basic_api/03_reg_vector_compute/truncate/truncate.asc+5-0
@@ -9,6 +9,11 @@
9*/9*/
10 10 
11 11 
12+/* !
13+ * \file truncate.asc
14+ * \brief 演示使用Trunc API实现向量截断取整计算功能，对输入向量逐元素截断取整并输出结果
15+ */
16+ 
12#include "acl/acl.h"17#include "acl/acl.h"
13#include "kernel_operator.h"18#include "kernel_operator.h"
14#include "data_utils.h"19#include "data_utils.h"
Mexamples/01_simd_cpp_api/02_features/03_basic_api/06_sync_control/reg_sync/CMakeLists.txt+3-16
@@ -11,13 +11,8 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14-if(DEFINED RUN_MODE)14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15- set(CMAKE_ASC_RUN_MODE ${RUN_MODE} CACHE STRING "Alias of RUN_MODE")15+set(CMAKE_ASC_ARCHITECTURES "dav-3510" CACHE STRING "NPU architecture: dav-3510")
16-endif()
17- 
18-if(DEFINED NPU_ARCH)
19- set(CMAKE_ASC_ARCHITECTURES ${NPU_ARCH} CACHE STRING "Alias of NPU_ARCH")
20-endif()
21 16 
22find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
23 18 
@@ -27,14 +22,6 @@ add_executable(demo
27 reg_sync.asc22 reg_sync.asc
28)23)
29 24 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36-set(NPU_ARCH "dav-3510" CACHE STRING "NPU ARCH, e.g. dav-3510")
37- 
38message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}")25message(STATUS "[INFO] SCENARIO_NUM is ${SCENARIO_NUM}")
39 26 
40if(NOT DEFINED SCENARIO_NUM)27if(NOT DEFINED SCENARIO_NUM)
@@ -50,6 +37,6 @@ target_compile_definitions(demo PRIVATE
50)37)
51 38 
52target_compile_options(demo PRIVATE39target_compile_options(demo PRIVATE
53- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${NPU_ARCH}>40+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
54 $<$<COMPILE_LANGUAGE:ASC>:--cce-simd-vf-fusion=false>41 $<$<COMPILE_LANGUAGE:ASC>:--cce-simd-vf-fusion=false>
55)42)
Mexamples/01_simd_cpp_api/02_features/03_basic_api/06_sync_control/reg_sync/README.md+9-9
@@ -76,19 +76,19 @@
76- 样例执行76- 样例执行
77 ```bash77 ```bash
78 SCENARIO=178 SCENARIO=1
79- mkdir -p build && cd build; # 创建并进入build目录79+ mkdir -p build && cd build; # 创建并进入build目录
80- cmake -DNPU_ARCH=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # 编译工程（默认npu模式）80+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # 编译工程（默认npu模式）
81- python3 ../scripts/gen_data.py -scenarioNum=$SCENARIO # 生成测试输入数据81+ python3 ../scripts/gen_data.py -scenarioNum=$SCENARIO # 生成测试输入数据
82- ./demo # 执行编译生成的可执行程序，执行样例82+ ./demo # 执行编译生成的可执行程序，执行样例
83 ```83 ```
84 84 
85- 使用 CPU调试 或 NPU仿真 模式时，添加 `-DRUN_MODE=cpu` 或 `-DRUN_MODE=sim` 参数即可。85+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
86 86 
87 示例如下：87 示例如下：
88 ```bash88 ```bash
89 SCENARIO=189 SCENARIO=1
90- cmake -DRUN_MODE=cpu -DNPU_ARCH=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # cpu调试模式90+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # cpu调试模式
91- cmake -DRUN_MODE=sim -DNPU_ARCH=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # NPU仿真模式91+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 -DSCENARIO_NUM=$SCENARIO ..;make -j; # NPU仿真模式
92 ```92 ```
93 93 
94 > **注意：** 切换编译模式或场景前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。94 > **注意：** 切换编译模式或场景前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
@@ -97,8 +97,8 @@
97 97 
98| 选项　　　　　 | 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|98| 选项　　　　　 | 可选值　　　　　　　　　　　| 说明　　　　　　　　　　　　　　　　　　　　　　　|
99| ----------------| -----------------------------| ---------------------------------------------------|99| ----------------| -----------------------------| ---------------------------------------------------|
100-| `RUN_MODE`　　 | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|100+| `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真　　　　　　　|
101-| `NPU_ARCH`　　 | `dav-3510`　　　　　　　　　| NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |101+| `CMAKE_ASC_ARCHITECTURES` | `dav-3510` | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
102| `SCENARIO_NUM` | `1`、`2`　　　　　　　　　　| 场景编号：1=读写依赖，2=写写依赖　　　　　　　　　|102| `SCENARIO_NUM` | `1`、`2`　　　　　　　　　　| 场景编号：1=读写依赖，2=写写依赖　　　　　　　　　|
103 103 
104- 执行结果104- 执行结果
Mexamples/01_simd_cpp_api/02_features/03_basic_api/06_sync_control/reg_sync/reg_sync.asc+1-0
@@ -108,6 +108,7 @@ __global__ __vector__ void reg_sync(GM_ADDR x, GM_ADDR y, GM_ADDR z)
108 Kernel<float> kernel;108 Kernel<float> kernel;
109 kernel.Init(x, y, z, 1024);109 kernel.Init(x, y, z, 1024);
110 kernel.Process();110 kernel.Process();
111+ AscendC::PipeBarrier<PIPE_ALL>();
111}112}
112 113 
113static bool CompareResult(const void* outputData, uint32_t outSize)114static bool CompareResult(const void* outputData, uint32_t outSize)