Star686
998
代码
代码
Issues99
Pull Requests189
Actions
项目讨论
Wiki
分析
Star686
998
已合并
高阶API Math样例整改 #1429
lipschitz_von创建于 4月3日
高阶API Math样例整改 #1429
已合并
lipschitz_von创建于 4月3日
从lipschitz_von/asc-devkitreorganizer-next-highorder-api-math01合入到cann/asc-devkitmaster
共 37 个文件变更+1188-604
Mexamples/01_simd_cpp_api/03_libraries/12_math/acosh/CMakeLists.txt+6-9
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -25,15 +28,9 @@ target_link_libraries(demo PRIVATE
25 platform28 platform
26 m29 m
27 dl30 dl
31+ graph_base
28)32)
29 33 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE34target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>35+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38- # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>36+)
39-)
Mexamples/01_simd_cpp_api/03_libraries/12_math/acosh/README.md+66-24
@@ -2,7 +2,19 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Acosh高阶API的算子实现。样例按元素做双曲反余弦函数计算。5+本样例基于Acosh高阶API计算反双曲余弦函数。
D
DDragonBornHD844月22日

其他接口提示？

likedislike
6+ 
7+> **涉及样例：** 除本样例使用的 `Acosh` 接口外，Ascend C 还提供了以下三角函数相关高阶API接口，除sincos外实现方式基本一致，如需调用替换接口名即可：
8+>
9+> - **acos**：反余弦函数。
10+> - **asin**：反正弦函数。
11+> - **asinh**：反双曲正弦函数。
12+> - **atanh**：反双曲正切函数。
13+> - **cos**：余弦函数。
14+> - **cosh**：双曲余弦函数。
15+> - **sinh**：双曲正弦函数。
16+> - **tan**：正切函数。
17+> - **sincos**：正弦余弦函数，分别计算正弦和余弦，调用时需要两个输出Tensor。
6 18 
7## 支持的产品19## 支持的产品
8 20 
@@ -12,69 +24,99 @@
12 24 
13## 目录结构介绍25## 目录结构介绍
14 26 
15-```27+```plain
16├── acosh28├── acosh
17│ ├── CMakeLists.txt // 编译工程文件29│ ├── CMakeLists.txt // 编译工程文件
18-│ └── acosh.asc // Ascend C算子实现 & 调用样例30+│ └── acosh.asc // Ascend C样例实现 & 调用样例
19```31```
20 32 
21-## 算子描述33+## 样例描述
22 34 
23-- 算子功能： 35+- 样例功能：
24- 按元素做双曲反余弦函数计算，计算公式如下： 36+ 按元素做双曲反余弦函数计算，计算公式如下：
25 $$dstTensor_i = Acosh(srcTensor_i)$$37 $$dstTensor_i = Acosh(srcTensor_i)$$
26 $$Acosh(x)=\begin{cases}Nan, & x < 1 \\ \ln(x+\sqrt{x^{2}-1}), & x > 1\end{cases}$$38 $$Acosh(x)=\begin{cases}Nan, & x < 1 \\ \ln(x+\sqrt{x^{2}-1}), & x > 1\end{cases}$$
27-- 算子规格： 39+- 样例规格：
28 <table>40 <table>
29- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> acosh </td></tr>41+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> acosh </td></tr>
30 42 
31- <tr><td rowspan="3" align="center">算子输入</td></tr>43+ <tr><td rowspan="3" align="center">样例输入</td></tr>
32 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>44 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
33- <tr><td align="center">src</td><td align="center">16</td><td align="center">float</td><td align="center">ND</td></tr>45+ <tr><td align="center">src</td><td align="center">[1, 16]</td><td align="center">float</td><td align="center">ND</td></tr>
34- <tr><td rowspan="2" align="center">算子输出</td></tr>46+ <tr><td rowspan="2" align="center">样例输出</td></tr>
35- <tr><td align="center">dst</td><td align="center">16</td><td align="center">float</td><td align="center">ND</td></tr>47+ <tr><td align="center">dst</td><td align="center">[1, 16]</td><td align="center">float</td><td align="center">ND</td></tr>
36 48 
37 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">acosh_custom</td></tr>49 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">acosh_custom</td></tr>
38 </table>50 </table>
39 51 
40-- 算子实现： 52+- 样例实现：
41- 本样例中实现的是固定shape为输入src[16]，输出dst[16]的acosh_custom算子。53+ 本样例中实现的是固定shape为输入src[1, 16]，输出dst[1, 16]的acosh_custom。
42 54 
43- - Kernel实现 55+ - Kernel实现
44- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Acosh高阶API接口完成Acosh计算，得到最终结果，再搬出到外部存储上。
45 56 
46- acosh_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，Compute任务负责对srcLocal执行Acosh计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。57+ 使用Acosh高阶API接口完成反双曲余弦计算。
47 58 
48- - 调用实现 59+ - Tiling实现
60+ 
61+ Host侧通过GetAcoshMaxMinTmpSize获取Acosh接口计算所需的最大和最小临时空间。
62+ 
63+ - 调用实现
49 使用内核调用符<<<>>>调用核函数。64 使用内核调用符<<<>>>调用核函数。
50 65 
51## 编译运行 66## 编译运行
52 67 
53-在本样例根目录下执行如下步骤，编译并执行算子。68+在本样例根目录下执行如下步骤，编译并执行样例。
69+ 
54- 配置环境变量 70- 配置环境变量
55 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。71 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
56 - 默认路径，root用户安装CANN软件包72 - 默认路径，root用户安装CANN软件包
73+ 
57 ```bash74 ```bash
58 source /usr/local/Ascend/cann/set_env.sh75 source /usr/local/Ascend/cann/set_env.sh
59 ```76 ```
60 77 
61 - 默认路径，非root用户安装CANN软件包78 - 默认路径，非root用户安装CANN软件包
79+ 
62 ```bash80 ```bash
63 source $HOME/Ascend/cann/set_env.sh81 source $HOME/Ascend/cann/set_env.sh
64 ```82 ```
65 83 
66 - 指定路径install_path，安装CANN软件包84 - 指定路径install_path，安装CANN软件包
85+ 
67 ```bash86 ```bash
68 source ${install_path}/cann/set_env.sh87 source ${install_path}/cann/set_env.sh
69 ```88 ```
70- 89+ 
71- 样例执行90- 样例执行
91+ 
72 ```bash92 ```bash
73- mkdir -p build && cd build; # 创建并进入build目录93+ mkdir -p build && cd build; # 创建并进入build目录
74- cmake ..;make -j; # 编译工程94+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式
75- ./demo # 执行编译生成的可执行程序，执行样例95+ ./demo # 执行编译生成的可执行程序，执行样例
76 ```96 ```
97+ 
98+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
99+
100+ 示例如下：
101+ 
D
DDragonBornHD844月25日

多余的空行

likedislike
102+ ```bash
103+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式
104+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式
105+ ```
106+ 
107+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
108+ 
109+- 编译选项说明
110+ 
111+ | 选项 | 可选值 | 说明 |
112+ |------|--------|------|
113+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
114+ | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2/A3 系列，dav-3510 对应 Ascend 950PR/Ascend 950DT |
115+ 
116+- 执行结果
117+ 
77 执行结果如下，说明精度对比成功。118 执行结果如下，说明精度对比成功。
119+ 
78 ```bash120 ```bash
79 test pass!121 test pass!
80- ```122+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/acosh/acosh.asc+28-10
@@ -11,13 +11,22 @@
11 11 
12/* !12/* !
13 * \file acosh.asc13 * \file acosh.asc
14- * \brief14+ * \brief Acosh样例实现，计算反双曲余弦函数
15 */15 */
16 16 
17#include <random>17#include <random>
18#include "acl/acl.h"18#include "acl/acl.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20+#include "tiling/tiling_api.h"
D
DDragonBornHD844月22日

代码几乎没注释

likedislike
20 21 
22+#ifdef ASCENDC_CPU_DEBUG
23+#include "cpu_debug_launch.h"
24+#endif
25+ 
26+/**
27+ * @brief Acosh核函数类，实现反双曲余弦函数计算
28+ * @tparam T 输入输出数据类型
29+ */
21template <typename T>30template <typename T>
22class KernelAcosh {31class KernelAcosh {
23public:32public:
@@ -39,13 +48,11 @@ public:
39 pipe->InitBuffer(buf, tmpBufSize * sizeof(T));48 pipe->InitBuffer(buf, tmpBufSize * sizeof(T));
40 }49 }
41 }50 }
42- __aicore__ inline void Process(uint32_t tmpBufSize)51+ __aicore__ inline void Process()
43 {52 {
44- AscendC::AscendCUtils::SetOverflow(1);
45 CopyIn();53 CopyIn();
46 Compute();54 Compute();
47 CopyOut();55 CopyOut();
48- AscendC::AscendCUtils::SetOverflow(0);
49 }56 }
50 57 
51 __aicore__ inline void CopyIn()58 __aicore__ inline void CopyIn()
@@ -64,6 +71,14 @@ public:
64 if (tmpBufSize > 0) {71 if (tmpBufSize > 0) {
65 temp = buf.Get<uint8_t>();72 temp = buf.Get<uint8_t>();
66 }73 }
74+ // 使用Acosh高阶API计算反双曲余弦函数
75+ // 模板参数：
76+ // - T: 输入输出数据类型
77+ // 参数说明：
78+ // - dstLocal: 输出Tensor，存储计算结果
79+ // - srcLocal: 输入Tensor
80+ // - temp: 临时空间（可选）
81+ // - calCount: 计算元素个数（可选）
67 if ((tmpBufSize > 0) && calCount > 0) {82 if ((tmpBufSize > 0) && calCount > 0) {
68 AscendC::Acosh<T>(dstLocal, srcLocal, temp, calCount);83 AscendC::Acosh<T>(dstLocal, srcLocal, temp, calCount);
69 } else if (tmpBufSize > 0) {84 } else if (tmpBufSize > 0) {
@@ -95,15 +110,14 @@ private:
95 uint32_t tmpBufSize = 0;110 uint32_t tmpBufSize = 0;
96};111};
97 112 
98-__global__ __vector__ void acosh_custom(GM_ADDR srcGm, GM_ADDR dstGm)113+__global__ __vector__ void acosh_custom(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t tmpBufSize)
99{114{
100 AscendC::TPipe pipe;115 AscendC::TPipe pipe;
101 constexpr uint32_t srcSize = 16;116 constexpr uint32_t srcSize = 16;
102- constexpr uint32_t tmpBufSize = 0;
103 constexpr uint32_t calCount = 16;117 constexpr uint32_t calCount = 16;
104 KernelAcosh<float> op;118 KernelAcosh<float> op;
105 op.Init(srcGm, dstGm, srcSize, tmpBufSize, calCount, &pipe);119 op.Init(srcGm, dstGm, srcSize, tmpBufSize, calCount, &pipe);
106- op.Process(tmpBufSize);120+ op.Process();
107}121}
108 122 
109static bool CompareResult(const void* outputData, const void* goldenData, uint32_t outSize)123static bool CompareResult(const void* outputData, const void* goldenData, uint32_t outSize)
@@ -160,6 +174,11 @@ int32_t main(int32_t argc, char* argv[])
160 size_t param2FileSize = 16 * sizeof(float);174 size_t param2FileSize = 16 * sizeof(float);
161 uint32_t numBlocks = 1;175 uint32_t numBlocks = 1;
162 176 
177+ ge::Shape shape{{16}};
178+ uint32_t maxValue = 0;
179+ uint32_t minValue = 0;
180+ AscendC::GetAcoshMaxMinTmpSize(shape, sizeof(float), false, maxValue, minValue);
181+ 
163 auto genData = gen_golden_data();182 auto genData = gen_golden_data();
164 auto src = genData.src;183 auto src = genData.src;
165 auto golden = genData.golden;184 auto golden = genData.golden;
@@ -183,11 +202,10 @@ int32_t main(int32_t argc, char* argv[])
183 aclrtMallocHost((void**)(&param2Host), param2FileSize);202 aclrtMallocHost((void**)(&param2Host), param2FileSize);
184 aclrtMalloc((void**)&param2Device, param2FileSize, ACL_MEM_MALLOC_HUGE_FIRST);203 aclrtMalloc((void**)&param2Device, param2FileSize, ACL_MEM_MALLOC_HUGE_FIRST);
185 204 
186- acosh_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device);205+ acosh_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, minValue);
187 aclrtSynchronizeStream(stream);206 aclrtSynchronizeStream(stream);
188 207 
189 aclrtFree(param1Device);208 aclrtFree(param1Device);
190- aclrtFreeHost(param1Host);
191 aclrtMemcpy(param2Host, param2FileSize, param2Device, param2FileSize, ACL_MEMCPY_DEVICE_TO_HOST);209 aclrtMemcpy(param2Host, param2FileSize, param2Device, param2FileSize, ACL_MEMCPY_DEVICE_TO_HOST);
192 210 
193 bool goldenResult = true;211 bool goldenResult = true;
@@ -207,4 +225,4 @@ int32_t main(int32_t argc, char* argv[])
207 aclFinalize();225 aclFinalize();
208 226 
209 return 0;227 return 0;
210-}228+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/axpy_half_float/CMakeLists.txt+6-9
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -25,15 +28,9 @@ target_link_libraries(demo PRIVATE
25 platform28 platform
26 m29 m
27 dl30 dl
31+ graph_base
28)32)
29 33 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE34target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>35+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38- # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>36+)
39-)
Mexamples/01_simd_cpp_api/03_libraries/12_math/axpy_half_float/README.md+61-27
@@ -1,15 +1,18 @@
1-# Axpy算子直调样例1+# Axpy样例
2+ 
2## 概述3## 概述
3-本样例基于Axpy实现源操作数src中每个元素与标量求积后和目的操作数dst中的对应元素相加的功能。Axpy接口的源操作数和目的操作数的数据类型只能取三种组合：(half, half)、(float, float)、(half, float)。本样例中输入tensor和标量的数据类型为half，输出tensor数据类型为float。4+ 
4-本样例通过Ascend C编程语言实现了Axpy算子，使用<<<>>>内核调用符来完成算子核函数在NPU侧运行验证的基础流程，给出了对应的端到端实现。5+本样例基于Axpy高阶API实现源操作数src中每个元素与标量求积后和目的操作数dst中的对应元素相加的功能。Axpy接口的源操作数和目的操作数的数据类型只能取三种组合：(half, half)、(float, float)、(half, float)。本样例中输入tensor和标量的数据类型为half，输出tensor数据类型为float。
5 6 
6## 支持的产品7## 支持的产品
8+ 
7- Ascend 950PR/Ascend 950DT9- Ascend 950PR/Ascend 950DT
8- Atlas A3 训练系列产品/Atlas A3 推理系列产品10- Atlas A3 训练系列产品/Atlas A3 推理系列产品
9- Atlas A2 训练系列产品/Atlas A2 推理系列产品11- Atlas A2 训练系列产品/Atlas A2 推理系列产品
10 12 
11## 目录结构介绍13## 目录结构介绍
12-```14+ 
15+```plain
13├── axpy_half_float16├── axpy_half_float
14│ ├── scripts17│ ├── scripts
15│ │ ├── gen_data.py // 输入数据和真值数据生成脚本18│ │ ├── gen_data.py // 输入数据和真值数据生成脚本
@@ -19,62 +22,93 @@
19│ └── axpy_half_float.asc // Ascend C算子实现 & 调用样例22│ └── axpy_half_float.asc // Ascend C算子实现 & 调用样例
20```23```
21 24 
22-## 算子描述25+## 样例描述
23-- 算子功能： 26+ 
24- Axpy算子实现了源操作数src中每个元素与标量求积后和目的操作数dst中的对应元素相加，并返回计算结果的功能。27+- 样例功能：
28+ Axpy样例实现了源操作数src中每个元素与标量求积后和目的操作数dst中的对应元素相加，并返回计算结果的功能。
25 29 
26 对应的数学表达式为： 30 对应的数学表达式为：
27- ```31+ 
32+ $$
28 out = x * scalar + out33 out = x * scalar + out
29- ```34+ $$
30-- 算子规格： 35+ 
36+- 样例规格：
31 <table>37 <table>
32- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="5" align="center"> Axpy </td></tr>38+ <caption>表1：样例规格</caption>
33- <tr><td rowspan="2" align="center">算子输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td><td align="center">default</td></tr> <tr><td align="center">x</td><td align="center">4 * 128</td><td align="center">float16</td><td align="center">ND</td><td align="center">\</td></tr> <tr><td rowspan="1" align="center">算子输出</td><td align="center">out</td><td align="center">4 * 128</td><td align="center">float32</td><td align="center">ND</td><td align="center">\</td></tr> <tr><td rowspan="1" align="center">核函数名</td><td colspan="5" align="center">kernel_vec_ternary_scalar_Axpy_half_2_float</td></tr> 39+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="5" align="center"> Axpy </td></tr>
40+ <tr><td rowspan="2" align="center">样例输入</td><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td><td align="center">default</td></tr> <tr><td align="center">x</td><td align="center">[4, 128]</td><td align="center">float16</td><td align="center">ND</td><td align="center">\</td></tr> <tr><td rowspan="1" align="center">样例输出</td><td align="center">out</td><td align="center">[4, 128]</td><td align="center">float32</td><td align="center">ND</td><td align="center">\</td></tr> <tr><td rowspan="1" align="center">核函数名</td><td colspan="5" align="center">kernel_vec_ternary_scalar_Axpy_half_2_float</td></tr>
34 </table>41 </table>
35 42 
36-- 算子实现： 43+- 样例实现：
37- 本样例中实现的是固定shape为4*128的Axpy算子。44+ 本样例中实现的是固定shape为输入x[4, 128], 输出out[4, 128]的Axpy样例。
38- - Kernel实现 45+ - Kernel实现
39- Axpy算子的数学表达式为：
40- ```
41- out = x * scalar + out
42- ```
43- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用计算接口完成x乘以标量scalar再加上out中的原始值，得到最终结果，再搬出到外部存储上。
44 46 
45- Axpy算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor xGm搬运至Local Memory，存储在xLocal，Compute任务负责对xLocal执行相关操作，计算结果存储在outLocal中，CopyOut任务负责将输出数据从outLocal搬运至Global Memory上的输出Tensor outGm中。47+ 首先使用Duplicate接口将输出tensor初始化为0，然后使用Axpy接口完成x乘以标量scalar再加上out中的原始值，得到最终结果，再搬出到外部存储上。
48+ 
49+ - Tiling实现
50+ 
51+ Host侧通过GetAxpyMaxMinTmpSize获取Axpy接口计算所需的最大和最小临时空间。
46 52 
47 - 调用实现 53 - 调用实现
48 使用内核调用符<<<>>>调用核函数。54 使用内核调用符<<<>>>调用核函数。
49 55 
50## 编译运行 56## 编译运行
51-在本样例根目录下执行如下步骤，编译并执行算子。57+ 
58+在本样例根目录下执行如下步骤，编译并执行样例。
59+ 
52- 配置环境变量 60- 配置环境变量
53 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。61 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
54 - 默认路径，root用户安装CANN软件包62 - 默认路径，root用户安装CANN软件包
63+ 
55 ```bash64 ```bash
56 source /usr/local/Ascend/cann/set_env.sh65 source /usr/local/Ascend/cann/set_env.sh
57 ```66 ```
58- 67+ 
59 - 默认路径，非root用户安装CANN软件包68 - 默认路径，非root用户安装CANN软件包
69+ 
60 ```bash70 ```bash
61 source $HOME/Ascend/cann/set_env.sh71 source $HOME/Ascend/cann/set_env.sh
62 ```72 ```
63 73 
64 - 指定路径install_path，安装CANN软件包74 - 指定路径install_path，安装CANN软件包
75+ 
65 ```bash76 ```bash
66 source ${install_path}/cann/set_env.sh77 source ${install_path}/cann/set_env.sh
67 ```78 ```
68 79 
69- 样例执行80- 样例执行
81+ 
70 ```bash82 ```bash
71- mkdir -p build && cd build; # 创建并进入build目录83+ mkdir -p build && cd build; # 创建并进入build目录
72- cmake ..;make -j; # 编译工程84+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式
73 python3 ../scripts/gen_data.py # 生成测试输入数据85 python3 ../scripts/gen_data.py # 生成测试输入数据
74- ./demo # 执行编译生成的可执行程序，执行样例86+ ./demo # 执行编译生成的可执行程序，执行样例
75 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确87 python3 ../scripts/verify_result.py output/output.bin output/golden.bin # 验证输出结果是否正确，确认算法逻辑正确
76 ```88 ```
89+ 
90+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
91+
92+ 示例如下：
93+ 
94+ ```bash
95+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式
96+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式
97+ ```
98+ 
99+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
100+ 
101+- 编译选项说明
102+ 
103+ | 选项 | 可选值 | 说明 |
104+ |------|--------|------|
105+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
106+ | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2/A3 系列，dav-3510 对应 Ascend 950PR/Ascend 950DT |
107+ 
108+- 执行结果
109+ 
77 执行结果如下，说明精度对比成功。110 执行结果如下，说明精度对比成功。
111+ 
78 ```bash112 ```bash
79 test pass!113 test pass!
80- ```114+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/axpy_half_float/axpy_half_float.asc+36-10
@@ -11,22 +11,32 @@
11 11 
12/* !12/* !
13 * \file axpy_half_float.asc13 * \file axpy_half_float.asc
14- * \brief14+ * \brief 基于Axpy高阶API实现half到float类型的向量标量乘加运算
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20+#include "tiling/tiling_api.h"
20 21 
22+#ifdef ASCENDC_CPU_DEBUG
D
DDragonBornHD844月16日

加一下tiling相关的逻辑，代码对应注释加一下

likedislike
23+#include "cpu_debug_launch.h"
24+#endif
25+ 
26+/**
27+ * @brief Axpy算子Kernel类，实现向量标量乘加运算
28+ */
21class KernelAxpy {29class KernelAxpy {
22public:30public:
23 __aicore__ inline KernelAxpy() {}31 __aicore__ inline KernelAxpy() {}
24- __aicore__ inline void Init(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm)32+ __aicore__ inline void Init(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm, uint32_t tmpBufSize, AscendC::TPipe* pipeIn)
25 {33 {
34+ pipe = pipeIn;
26 srcGlobal.SetGlobalBuffer((__gm__ half*)srcGm);35 srcGlobal.SetGlobalBuffer((__gm__ half*)srcGm);
27 dstGlobal.SetGlobalBuffer((__gm__ float*)dstGm);36 dstGlobal.SetGlobalBuffer((__gm__ float*)dstGm);
28- pipe.InitBuffer(outQueueDst, 1, 512 * sizeof(float));37+ pipe->InitBuffer(outQueueDst, 1, 512 * sizeof(float));
29- pipe.InitBuffer(inQueueSrc, 1, 512 * sizeof(half));38+ pipe->InitBuffer(inQueueSrc, 1, 512 * sizeof(half));
39+ pipe->InitBuffer(buf, tmpBufSize * sizeof(uint8_t));
30 }40 }
31 __aicore__ inline void Process()41 __aicore__ inline void Process()
32 {42 {
@@ -45,9 +55,18 @@ private:
45 {55 {
46 AscendC::LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();56 AscendC::LocalTensor<half> srcLocal = inQueueSrc.DeQue<half>();
47 AscendC::LocalTensor<float> dstLocal = outQueueDst.AllocTensor<float>();57 AscendC::LocalTensor<float> dstLocal = outQueueDst.AllocTensor<float>();
58+ AscendC::LocalTensor<uint8_t> sharedTmpBuffer = buf.Get<uint8_t>();
48 59
60+ // 初始化输出tensor为0
49 AscendC::Duplicate(dstLocal, 0.0f, 512);61 AscendC::Duplicate(dstLocal, 0.0f, 512);
50- AscendC::Axpy(dstLocal, srcLocal, (half)2.0, 64, 8, { 1, 1, 8, 4 });62+ // 使用Axpy高阶API进行向量标量乘加运算
63+ // 参数说明：
64+ // - dstLocal: 输出Tensor，存储计算结果（float类型）
65+ // - srcLocal: 输入Tensor（half类型）
66+ // - scalar: 标量值（half类型）
67+ // - sharedTmpBuffer: 临时空间，用于类型转换
68+ // - calCount: 计算元素个数
69+ AscendC::Axpy(dstLocal, srcLocal, static_cast<half>(2.0), sharedTmpBuffer, 512);
51 70
52 outQueueDst.EnQue<float>(dstLocal);71 outQueueDst.EnQue<float>(dstLocal);
53 inQueueSrc.FreeTensor(srcLocal);72 inQueueSrc.FreeTensor(srcLocal);
@@ -59,17 +78,19 @@ private:
59 outQueueDst.FreeTensor(dstLocal);78 outQueueDst.FreeTensor(dstLocal);
60 }79 }
61private:80private:
62- AscendC::TPipe pipe;81+ AscendC::TPipe* pipe;
63 AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueSrc;82 AscendC::TQue<AscendC::TPosition::VECIN, 1> inQueueSrc;
64 AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueDst;83 AscendC::TQue<AscendC::TPosition::VECOUT, 1> outQueueDst;
84+ AscendC::TBuf<AscendC::TPosition::VECCALC> buf;
65 AscendC::GlobalTensor<half> srcGlobal;85 AscendC::GlobalTensor<half> srcGlobal;
66 AscendC::GlobalTensor<float> dstGlobal;86 AscendC::GlobalTensor<float> dstGlobal;
67};87};
68 88 
69-__global__ __vector__ void kernel_vec_ternary_scalar_Axpy_half_2_float(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm)89+__global__ __vector__ void kernel_vec_ternary_scalar_Axpy_half_2_float(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm, uint32_t tmpBufSize)
70{90{
91+ AscendC::TPipe pipe;
71 KernelAxpy op;92 KernelAxpy op;
72- op.Init(srcGm, dstGm);93+ op.Init(srcGm, dstGm, tmpBufSize, &pipe);
73 op.Process();94 op.Process();
74}95}
75 96 
@@ -78,6 +99,11 @@ int32_t main(int32_t argc, char *argv[]) {
78 size_t inputByteSize = 4 * 128 * sizeof(uint16_t);99 size_t inputByteSize = 4 * 128 * sizeof(uint16_t);
79 size_t outputByteSize = 4 * 128 * sizeof(uint32_t);100 size_t outputByteSize = 4 * 128 * sizeof(uint32_t);
80 101 
102+ ge::Shape shape{{512}};
103+ uint32_t maxValue = 0;
104+ uint32_t minValue = 0;
105+ AscendC::GetAxpyMaxMinTmpSize(shape, sizeof(half), false, maxValue, minValue);
106+ 
81 int32_t deviceId = 0;107 int32_t deviceId = 0;
82 aclrtSetDevice(deviceId);108 aclrtSetDevice(deviceId);
83 aclrtStream stream = nullptr;109 aclrtStream stream = nullptr;
@@ -96,7 +122,7 @@ int32_t main(int32_t argc, char *argv[]) {
96 aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize,122 aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize,
97 ACL_MEMCPY_HOST_TO_DEVICE);123 ACL_MEMCPY_HOST_TO_DEVICE);
98 124 
99- kernel_vec_ternary_scalar_Axpy_half_2_float<<<numBlocks, nullptr, stream>>>(xDevice, outDevice);125+ kernel_vec_ternary_scalar_Axpy_half_2_float<<<numBlocks, nullptr, stream>>>(xDevice, outDevice, minValue);
100 aclrtSynchronizeStream(stream);126 aclrtSynchronizeStream(stream);
101 127 
102 aclrtMemcpy(outHost, outputByteSize, outDevice, outputByteSize,128 aclrtMemcpy(outHost, outputByteSize, outDevice, outputByteSize,
@@ -113,4 +139,4 @@ int32_t main(int32_t argc, char *argv[]) {
113 aclFinalize();139 aclFinalize();
114 140 
115 return 0;141 return 0;
116-}142+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/clamp/CMakeLists.txt+6-8
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 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-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -25,14 +28,9 @@ target_link_libraries(demo PRIVATE
25 platform28 platform
26 m29 m
27 dl30 dl
31+ graph_base
28)32)
29 33 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE34target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>35+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38-)36+)
Mexamples/01_simd_cpp_api/03_libraries/12_math/clamp/README.md+78-36
@@ -2,7 +2,8 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Clamp高阶API的算子实现。样例将输入中除nan值以外大于max的数替换为max，小于min的数替换为min，小于等于max和大于等于min的数保持不变，作为输出。当min大于max时，将除nan值外所有值替换为max。min和max可以为标量或LocalTensor。 5+本样例基于Clamp高阶API实现将输入中除nan值以外的数截断到区间[min, max]的功能。
6+当min大于max时，将除nan值外所有值替换为max。min和max均可以为标量或张量。
6 7 
7## 支持的产品8## 支持的产品
8 9 
@@ -10,87 +11,128 @@
10 11 
11## 目录结构介绍12## 目录结构介绍
12 13 
13-```14+```plain
14├── clamp15├── clamp
15│ ├── scripts16│ ├── scripts
16-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本17+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
17│ ├── CMakeLists.txt // 编译工程文件18│ ├── CMakeLists.txt // 编译工程文件
18│ ├── data_utils.h // 数据读入写出函数19│ ├── data_utils.h // 数据读入写出函数
19-│ └── clamp.asc // Ascend C算子实现 & 调用样例20+│ └── clamp.asc // Ascend C样例实现 & 调用样例
20```21```
21 22 
22-## 算子描述23+## 样例描述
24+ 
25+- 样例功能：
26+ 将输入中大于max的非NaN值替换为max，小于min的非NaN值替换为min，小于等于max和大于等于min的数保持不变，作为输出。当min大于max时，将所有非NaN值替换为max。min和max可以为标量或张量。
23 27 
24-- 算子功能：
25- 将输入中除nan值以外大于max的数替换为max，小于min的数替换为min，小于等于max和大于等于min的数保持不变，作为输出。当min大于max时，将除nan值外所有值替换为max。min和max可以为标量或LocalTensor。
26-
27 计算公式如下： 28 计算公式如下：
29+ 
28 $$30 $$
29 dst_i = Clamp(src_i, min_i, max_i)31 dst_i = Clamp(src_i, min_i, max_i)
30 $$32 $$
31-$$
32-dst_i =
33-\begin{cases}
34-min_i, & src_i < min_i \\
35-src_i, & min_i \le src_i \le max_i \\
36-max_i, & src_i > max_i \\
37-\end{cases}
38-$$
39 33 
40-- 算子规格： 34+ $$
35+ Clamp(src_i, min_i, max_i) =
36+ \begin{cases}
37+ min_i, & src_i < min_i \\
38+ src_i, & min_i \le src_i \le max_i \\
39+ max_i, & src_i > max_i \\
40+ \end{cases}
41+ $$
42+ 
43+- 样例规格：
41 <table>44 <table>
42- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> clamp </td></tr>45+ <caption>表1：样例规格</caption>
46+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> clamp </td></tr>
43 47 
44- <tr><td rowspan="5" align="center">算子输入</td></tr>48+ <tr><td rowspan="5" align="center">样例输入</td></tr>
45 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>49 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
46- <tr><td align="center">src</td><td align="center">128</td><td align="center">float</td><td align="center">ND</td></tr>50+ <tr><td align="center">src</td><td align="center">[1, 128]</td><td align="center">float</td><td align="center">ND</td></tr>
47- <tr><td align="center">src_min</td><td align="center">128</td><td align="center">float</td><td align="center">ND</td></tr>51+ <tr><td align="center">src_min</td><td align="center">[1, 128]</td><td align="center">float</td><td align="center">ND</td></tr>
48- <tr><td align="center">src_max</td><td align="center">128</td><td align="center">float</td><td align="center">ND</td></tr>52+ <tr><td align="center">src_max</td><td align="center">[1, 128]</td><td align="center">float</td><td align="center">ND</td></tr>
49- <tr><td rowspan="2" align="center">算子输出</td></tr>53+ <tr><td rowspan="2" align="center">样例输出</td></tr>
50- <tr><td align="center">dst</td><td align="center">128</td><td align="center">float</td><td align="center">ND</td></tr>54+ <tr><td align="center">dst</td><td align="center">[1, 128]</td><td align="center">float</td><td align="center">ND</td></tr>
51 55 
52 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">clamp_custom</td></tr>56 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">clamp_custom</td></tr>
53 </table>57 </table>
54 58 
55-- 算子实现： 59+- 场景说明：
56- 本样例中实现的是固定shape为输入src[128]、src_min[128]、src_max[128]，输出dst[128]的clamp_custom算子。60+ <table>
61+ <caption>表2：scalarType参数说明</caption>
62+ <tr><td align="center">scalarType</td><td align="center">min类型</td><td align="center">max类型</td><td align="center">说明</td></tr>
63+ <tr><td align="center">1</td><td align="center">张量</td><td align="center">张量</td><td align="center">min和max都是张量</td></tr>
64+ <tr><td align="center">2</td><td align="center">张量</td><td align="center">标量</td><td align="center">min是张量，max是标量</td></tr>
65+ <tr><td align="center">3</td><td align="center">标量</td><td align="center">张量</td><td align="center">min是标量，max是张量</td></tr>
66+ <tr><td align="center">4</td><td align="center">标量</td><td align="center">标量</td><td align="center">min和max都是标量</td></tr>
67+ </table>
57 68 
58- - Kernel实现 69+- 样例实现：
59- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Clamp高阶API接口完成Clamp计算，得到最终结果，再搬出到外部存储上。
60 70 
61- clamp_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm、minGm、maxGm存储在srcLocal、minLocal、maxLocal中，Compute任务负责对srcLocal、minLocal、maxLocal执行Clamp计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。71+ 本样例中实现的是shape为输入src[128]、src_min[128]、src_max[128]，输出dst[128]的clamp_custom样例，支持min和max为张量或标量的4种场景组合。
D
DDragonBornHD844月22日

换行失效

likedislike
72+ 
73+ - Kernel实现
74+ 
75+ 使用Clamp高阶API接口完成Clamp计算，得到最终结果，再搬出到外部存储上。
76+ 
77+ - 调用实现
62 78 
63- - 调用实现
64 使用内核调用符<<<>>>调用核函数。79 使用内核调用符<<<>>>调用核函数。
65 80 
66## 编译运行 81## 编译运行
67 82 
68-在本样例根目录下执行如下步骤，编译并执行算子。83+在本样例根目录下执行如下步骤，编译并执行样例。
84+ 
69- 配置环境变量 85- 配置环境变量
70 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。86 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
71 - 默认路径，root用户安装CANN软件包87 - 默认路径，root用户安装CANN软件包
88+ 
72 ```bash89 ```bash
73 source /usr/local/Ascend/cann/set_env.sh90 source /usr/local/Ascend/cann/set_env.sh
74 ```91 ```
75 92 
76 - 默认路径，非root用户安装CANN软件包93 - 默认路径，非root用户安装CANN软件包
94+ 
77 ```bash95 ```bash
78 source $HOME/Ascend/cann/set_env.sh96 source $HOME/Ascend/cann/set_env.sh
79 ```97 ```
80 98 
81 - 指定路径install_path，安装CANN软件包99 - 指定路径install_path，安装CANN软件包
100+ 
82 ```bash101 ```bash
83 source ${install_path}/cann/set_env.sh102 source ${install_path}/cann/set_env.sh
84 ```103 ```
85- 104+ 
86- 样例执行105- 样例执行
106+ 
87 ```bash107 ```bash
88- mkdir -p build && cd build; # 创建并进入build目录108+ mkdir -p build && cd build; # 创建并进入build目录
89- cmake ..;make -j; # 编译工程109+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程，默认npu模式
90 python3 ../scripts/gen_data.py # 生成测试输入数据110 python3 ../scripts/gen_data.py # 生成测试输入数据
91- ./demo # 执行编译生成的可执行程序，执行样例111+ ./demo # 执行编译生成的可执行程序，执行样例
92 ```112 ```
113+ 
114+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
115+
116+ 示例如下：
117+ 
118+ ```bash
119+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
120+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
121+ ```
122+ 
123+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
124+ 
125+- 编译选项说明
126+ 
127+ | 选项 | 可选值 | 说明 |
128+ |------|--------|------|
129+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
130+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认） | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
131+ 
132+- 执行结果
133+ 
93 执行结果如下，说明精度对比成功。134 执行结果如下，说明精度对比成功。
135+ 
94 ```bash136 ```bash
95 test pass!137 test pass!
96- ```138+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/clamp/clamp.asc+23-4
@@ -11,13 +11,17 @@
11 11 
12/* !12/* !
13 * \file clamp.asc13 * \file clamp.asc
14- * \brief14+ * \brief Clamp样例实现，支持min和max为张量或标量的4种场景组合
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20 20 
21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
21enum ScalarType {25enum ScalarType {
22 BOTH_TENSOR = 1,26 BOTH_TENSOR = 1,
23 TENSOR_SCALAR,27 TENSOR_SCALAR,
@@ -25,6 +29,10 @@ enum ScalarType {
25 BOTN_SCALAR29 BOTN_SCALAR
26};30};
27 31 
32+/**
33+ * @brief Clamp核函数类，实现将输入值截断到[min, max]区间的功能
34+ * @tparam T 输入输出数据类型
35+ */
28template <typename T>36template <typename T>
29class ClampByMinMax {37class ClampByMinMax {
每天都要吃馒头
每天都要吃馒头每天都要吃馒头4月17日

Compute部分可以添加一些必要的注释

likedislike
30public:38public:
@@ -49,11 +57,9 @@ public:
49 }57 }
50 __aicore__ inline void Process()58 __aicore__ inline void Process()
51 {59 {
52- AscendC::AscendCUtils::SetOverflow(1);
53 CopyIn();60 CopyIn();
54 Compute();61 Compute();
55 CopyOut();62 CopyOut();
56- AscendC::AscendCUtils::SetOverflow(0);
57 }63 }
58 __aicore__ inline void CopyIn()64 __aicore__ inline void CopyIn()
59 {65 {
@@ -76,9 +82,20 @@ public:
76 AscendC::LocalTensor<T> maxLocal = inQueueMax.DeQue<T>();82 AscendC::LocalTensor<T> maxLocal = inQueueMax.DeQue<T>();
77 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();83 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();
78 Duplicate(dstLocal, (T)0, dataSize);84 Duplicate(dstLocal, (T)0, dataSize);
85+ // 使用Clamp接口将输入截断到[min, max]区间
86+ // 模板参数：
87+ // - T: 输入输出数据类型
88+ // 参数说明：
89+ // - dstLocal: 输出Tensor，存储计算结果
90+ // - srcLocal: 输入Tensor
91+ // - minLocal/minValue: 最小值（张量或标量）
92+ // - maxLocal/maxValue: 最大值（张量或标量）
93+ // - count: 计算元素个数
79 if (scalarType == ScalarType::BOTH_TENSOR) {94 if (scalarType == ScalarType::BOTH_TENSOR) {
D
DDragonBornHD844月22日

这里不同分支做的啥事儿也应该有注释

likedislike
95+ // 截断区间上下界均为张量
80 AscendC::Clamp(dstLocal, srcLocal, minLocal, maxLocal, count);96 AscendC::Clamp(dstLocal, srcLocal, minLocal, maxLocal, count);
81 } else if (scalarType == ScalarType::TENSOR_SCALAR) {97 } else if (scalarType == ScalarType::TENSOR_SCALAR) {
98+ // 截断区间下界为张量，上界为标量
82 event_t eventIdMte2ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_S));99 event_t eventIdMte2ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_S));
83 AscendC::SetFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);100 AscendC::SetFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);
84 AscendC::WaitFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);101 AscendC::WaitFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);
@@ -88,6 +105,7 @@ public:
88 AscendC::WaitFlag<AscendC::HardEvent::S_V>(eventIdSToV);105 AscendC::WaitFlag<AscendC::HardEvent::S_V>(eventIdSToV);
89 AscendC::Clamp(dstLocal, srcLocal, minLocal, maxValue, count);106 AscendC::Clamp(dstLocal, srcLocal, minLocal, maxValue, count);
90 } else if (scalarType == ScalarType::SCALAR_TENSOR) {107 } else if (scalarType == ScalarType::SCALAR_TENSOR) {
108+ // 截断区间下界为标量，上界为张量
91 event_t eventIdMte2ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_S));109 event_t eventIdMte2ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_S));
92 AscendC::SetFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);110 AscendC::SetFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);
93 AscendC::WaitFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);111 AscendC::WaitFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);
@@ -97,6 +115,7 @@ public:
97 AscendC::WaitFlag<AscendC::HardEvent::S_V>(eventIdSToV);115 AscendC::WaitFlag<AscendC::HardEvent::S_V>(eventIdSToV);
98 AscendC::Clamp(dstLocal, srcLocal, minValue, maxLocal, count);116 AscendC::Clamp(dstLocal, srcLocal, minValue, maxLocal, count);
99 } else {117 } else {
118+ // 截断区间上下界均为标量
100 event_t eventIdMte2ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_S));119 event_t eventIdMte2ToS = static_cast<event_t>(GetTPipePtr()->FetchEventID(AscendC::HardEvent::MTE2_S));
101 AscendC::SetFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);120 AscendC::SetFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);
102 AscendC::WaitFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);121 AscendC::WaitFlag<AscendC::HardEvent::MTE2_S>(eventIdMte2ToS);
@@ -262,4 +281,4 @@ int32_t main(int32_t argc, char* argv[])
262 aclFinalize();281 aclFinalize();
263 282 
264 return 0;283 return 0;
265-}284+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/cumsum/CMakeLists.txt+4-8
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -29,13 +32,6 @@ target_link_libraries(demo PRIVATE
29 graph_base32 graph_base
30)33)
31 34 
32-# ======================================================================================
33-# NPU 编译选项配置
34-#
35-# 说明：
36-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
37-# ======================================================================================
38target_compile_options(demo PRIVATE35target_compile_options(demo PRIVATE
39- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>36+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
40- # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>
41)37)
Mexamples/01_simd_cpp_api/03_libraries/12_math/cumsum/README.md+57-25
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例基于Kernel直调算子工程，介绍了调用CumSum高阶API实现cumsum单算子，用于对输入张量按行或列进行累加和操作，输出结果中每个元素都是输入张量中对应位置及之前所有行或列的元素累加和。5+本样例基于CumSum高阶API实现张量按行或列计算累加和的功能。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -11,41 +11,46 @@
11- Atlas A2 训练系列产品/Atlas A2 推理系列产品11- Atlas A2 训练系列产品/Atlas A2 推理系列产品
12 12 
13## 目录结构介绍13## 目录结构介绍
14-```14+ 
15+```plain
15├── cumsum16├── cumsum
16│ ├── scripts17│ ├── scripts
17-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本18+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
18│ ├── CMakeLists.txt // 编译工程文件19│ ├── CMakeLists.txt // 编译工程文件
19│ ├── data_utils.h // 数据读入写出函数20│ ├── data_utils.h // 数据读入写出函数
20-│ └── cumsum.asc // Ascend C算子实现 & 调用样例21+│ └── cumsum.asc // Ascend C样例实现 & 调用样例
21```22```
22 23 
23-## 算子描述24+## 样例描述
24-- 算子功能： 25+ 
25- cumsum单算子，用于对输入张量按行或列进行累加和操作，输出结果中每个元素都是输入张量中对应位置及之前所有行或列的元素累加和。26+- 样例功能：
26-- 算子规格： 27+ 对输入张量按行或列进行累加和操作，输出结果中每个元素都是输入张量中对应位置及之前所有行或列的元素累加和。
28+- 样例规格：
27 <table>29 <table>
28- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> cumsum </td></tr>30+ <caption>表1：样例规格</caption>
31+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> cumsum </td></tr>
29 32 
30- <tr><td rowspan="3" align="center">算子输入</td></tr>33+ <tr><td rowspan="3" align="center">样例输入</td></tr>
31 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>34 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
32- <tr><td align="center">src</td><td align="center">32 * 160</td><td align="center">float</td><td align="center">ND</td></tr>35+ <tr><td align="center">src</td><td align="center">[32, 160]</td><td align="center">float</td><td align="center">ND</td></tr>
33- 
34- <tr><td rowspan="3" align="center">算子输出</td></tr>
35- <tr><td align="center">dst</td><td align="center">32 * 160</td><td align="center">float</td><td align="center">ND</td></tr>
36- <tr><td align="center">lastRow</td><td align="center">160</td><td align="center">float</td><td align="center">ND</td></tr>
37 36 
37+ <tr><td rowspan="3" align="center">样例输出</td></tr>
38+ <tr><td align="center">dst</td><td align="center">[32, 160]</td><td align="center">float</td><td align="center">ND</td></tr>
39+ <tr><td align="center">lastRow</td><td align="center">[1, 160]</td><td align="center">float</td><td align="center">ND</td></tr>
38 40 
39 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">cumsum_custom</td></tr>41 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">cumsum_custom</td></tr>
40 </table>42 </table>
41 43 
42-- 算子实现： 44+- 样例实现：
43- 本样例中实现的是固定shape为输入src[32, 160]，输出dst[32, 160]、lastRow[160]的cumsum算子。45+ 本样例中实现的是shape为输入src[32, 160]，输出dst[32, 160]、lastRow[1, 160]的cumsum_custom样例。
44 46 
45- - Kernel实现 47+ - Kernel实现
46- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用CumSum高阶API接口完成cumsum计算，得到最终结果，再搬出到外部存储上。
47 48 
48- cumsum算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，Compute任务负责对srcLocal执行cumsum计算，计算结果存储在dstLocal、lastRowLocal中，CopyOut任务负责将输出数据从dstLocal、lastRowLocal搬运至Global Memory上的输出Tensor dstGm、lastRowGm。49+ 使用CumSum高阶API接口完成cumsum计算。
50+ 
51+ - Tiling实现
52+ 
53+ Host侧通过GetCumSumMaxMinTmpSize获取CumSum接口计算所需的最大和最小临时空间。
49 54 
50 - 调用实现 55 - 调用实现
51 使用内核调用符<<<>>>调用核函数。56 使用内核调用符<<<>>>调用核函数。
每天都要吃馒头
每天都要吃馒头每天都要吃馒头4月17日

不需要写 调用实现

likedislike
L
lipschitz_von
4月20日 评论：
@@ -53,31 +58,58 @@
53## 编译运行 58## 编译运行
54 59 
55在本样例根目录下执行如下步骤，编译并执行算子。60在本样例根目录下执行如下步骤，编译并执行算子。
61+ 
56- 配置环境变量 62- 配置环境变量
57 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。63 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
58 - 默认路径，root用户安装CANN软件包64 - 默认路径，root用户安装CANN软件包
65+ 
59 ```bash66 ```bash
60 source /usr/local/Ascend/cann/set_env.sh67 source /usr/local/Ascend/cann/set_env.sh
61 ```68 ```
62 69 
63 - 默认路径，非root用户安装CANN软件包70 - 默认路径，非root用户安装CANN软件包
71+ 
64 ```bash72 ```bash
65 source $HOME/Ascend/cann/set_env.sh73 source $HOME/Ascend/cann/set_env.sh
66 ```74 ```
67 75 
68 - 指定路径install_path，安装CANN软件包76 - 指定路径install_path，安装CANN软件包
77+ 
69 ```bash78 ```bash
70 source ${install_path}/cann/set_env.sh79 source ${install_path}/cann/set_env.sh
71 ```80 ```
72- 81+ 
73- 样例执行82- 样例执行
83+ 
74 ```bash84 ```bash
75- mkdir -p build && cd build; # 创建并进入build目录85+ mkdir -p build && cd build; # 创建并进入build目录
76- cmake ..;make -j; # 编译工程86+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式
77 python3 ../scripts/gen_data.py # 生成测试输入数据87 python3 ../scripts/gen_data.py # 生成测试输入数据
78- ./demo # 执行编译生成的可执行程序，执行样例88+ ./demo # 执行编译生成的可执行程序，执行样例
79 ```89 ```
90+ 
91+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
92+
93+ 示例如下：
94+ 
95+ ```bash
96+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式
97+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式
98+ ```
99+ 
100+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
101+ 
102+- 编译选项说明
103+ 
104+ | 选项 | 可选值 | 说明 |
105+ |------|--------|------|
106+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
107+ | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2/A3 系列，dav-3510 对应 Ascend 950PR/Ascend 950DT |
108+ 
109+- 执行结果
110+ 
80 执行结果如下，说明精度对比成功。111 执行结果如下，说明精度对比成功。
112+ 
81 ```bash113 ```bash
82 test pass!114 test pass!
83- ```115+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/cumsum/cumsum.asc+40-9
@@ -11,12 +11,17 @@
11 11 
12/* !12/* !
13 * \file cumsum.asc13 * \file cumsum.asc
14- * \brief14+ * \brief CumSum样例实现，对输入张量按行或列进行累加和操作
D
DDragonBornHD844月22日

没加上tmpsize的tiling函数

likedislike
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20+#include "tiling/tiling_api.h"
21+ 
22+#ifdef ASCENDC_CPU_DEBUG
23+#include "cpu_debug_launch.h"
24+#endif
20 25 
21constexpr int32_t BUFFER_NUM = 1;26constexpr int32_t BUFFER_NUM = 1;
22 27 
@@ -27,12 +32,17 @@ __aicore__ inline uint32_t Align32B(uint32_t len)
27 return (len + alignSize - 1) / alignSize * alignSize;32 return (len + alignSize - 1) / alignSize * alignSize;
28}33}
29 34 
35+/**
36+ * @brief CumSum核函数类，实现按行或列进行累加和操作
37+ * @tparam T 输入输出数据类型
38+ * @tparam CONFIG CumSum配置参数
39+ */
30template <typename T, const AscendC::CumSumConfig& CONFIG>40template <typename T, const AscendC::CumSumConfig& CONFIG>
31class KernelCumSum {41class KernelCumSum {
32public:42public:
33 __aicore__ inline KernelCumSum() {}43 __aicore__ inline KernelCumSum() {}
34 __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, GM_ADDR lastRowGm, uint32_t outter, uint32_t inner,44 __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, GM_ADDR lastRowGm, uint32_t outter, uint32_t inner,
35- uint32_t gmOutter, uint32_t gmInner, AscendC::TPipe* tPipe)45+ uint32_t gmOutter, uint32_t gmInner, uint32_t tmpBufSize, AscendC::TPipe* tPipe)
36 {46 {
37 this->outter = outter;47 this->outter = outter;
38 this->inner = inner;48 this->inner = inner;
@@ -47,6 +57,7 @@ public:
47 this->pipe->InitBuffer(inQueueX, BUFFER_NUM, Align32B<T>(outter * inner) * sizeof(T));57 this->pipe->InitBuffer(inQueueX, BUFFER_NUM, Align32B<T>(outter * inner) * sizeof(T));
48 this->pipe->InitBuffer(outQueue, BUFFER_NUM, Align32B<T>(outter * inner) * sizeof(T));58 this->pipe->InitBuffer(outQueue, BUFFER_NUM, Align32B<T>(outter * inner) * sizeof(T));
49 this->pipe->InitBuffer(lastRowQueue, BUFFER_NUM, Align32B<T>(inner) * sizeof(T));59 this->pipe->InitBuffer(lastRowQueue, BUFFER_NUM, Align32B<T>(inner) * sizeof(T));
60+ this->pipe->InitBuffer(buf, tmpBufSize * sizeof(uint8_t));
50 }61 }
51 __aicore__ inline void Process()62 __aicore__ inline void Process()
52 {63 {
@@ -72,8 +83,21 @@ private:
72 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();83 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();
73 AscendC::LocalTensor<T> lastRowLocal = lastRowQueue.AllocTensor<T>();84 AscendC::LocalTensor<T> lastRowLocal = lastRowQueue.AllocTensor<T>();
74 AscendC::LocalTensor<T> srcLocal = inQueueX.DeQue<T>();85 AscendC::LocalTensor<T> srcLocal = inQueueX.DeQue<T>();
86+ AscendC::LocalTensor<uint8_t> sharedTmpBuffer = buf.Get<uint8_t>();
87+ // outter：输入数据的外轴长度
88+ // inner： 输入数据的内轴长度
75 const AscendC::CumSumInfo cumSumInfo{outter, inner};89 const AscendC::CumSumInfo cumSumInfo{outter, inner};
D
DDragonBornHD844月25日

为啥这种关键代码什么注释都没有呢，这里cumSumInfo入参的outter inner不解释谁看得懂是干啥的？

likedislike
76- AscendC::CumSum<T, CONFIG>(dstLocal, lastRowLocal, srcLocal, cumSumInfo);90+ // 使用CumSum高阶API进行累加和计算
91+ // 模板参数：
92+ // - T: 输入输出数据类型
93+ // - CONFIG: CumSum配置参数，包含axis、isReverse等
94+ // 参数说明：
95+ // - dstLocal: 输出Tensor，存储累加结果
96+ // - lastRowLocal: 最后一行结果Tensor
97+ // - srcLocal: 输入Tensor
98+ // - sharedTmpBuffer: 临时空间
99+ // - cumSumInfo: 累加配置信息（outter, inner）
100+ AscendC::CumSum<T, CONFIG>(dstLocal, lastRowLocal, srcLocal, sharedTmpBuffer, cumSumInfo);
77 outQueue.EnQue<T>(dstLocal);101 outQueue.EnQue<T>(dstLocal);
78 lastRowQueue.EnQue<T>(lastRowLocal);102 lastRowQueue.EnQue<T>(lastRowLocal);
79 inQueueX.FreeTensor(srcLocal);103 inQueueX.FreeTensor(srcLocal);
@@ -99,6 +123,7 @@ private:
99 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX;123 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX;
100 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueue;124 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueue;
101 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> lastRowQueue;125 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> lastRowQueue;
126+ AscendC::TBuf<AscendC::QuePosition::VECCALC> buf;
102 127 
103 AscendC::TPipe* pipe;128 AscendC::TPipe* pipe;
104 129 
@@ -115,7 +140,7 @@ __aicore__ constexpr AscendC::CumSumConfig GetConfig()
115 140 
116constexpr AscendC::CumSumConfig CONFIG = GetConfig();141constexpr AscendC::CumSumConfig CONFIG = GetConfig();
117 142 
118-__global__ __vector__ void cumsum_custom(GM_ADDR srcGm, GM_ADDR dstGm, GM_ADDR lastRowGm, GM_ADDR tilingGm)143+__global__ __vector__ void cumsum_custom(GM_ADDR srcGm, GM_ADDR dstGm, GM_ADDR lastRowGm, GM_ADDR tilingGm, uint32_t tmpBufSize)
119{144{
120 AscendC::TPipe tPipe;145 AscendC::TPipe tPipe;
121 KernelCumSum<float, CONFIG> op;146 KernelCumSum<float, CONFIG> op;
@@ -123,7 +148,7 @@ __global__ __vector__ void cumsum_custom(GM_ADDR srcGm, GM_ADDR dstGm, GM_ADDR l
123 constexpr uint32_t inner = 160;148 constexpr uint32_t inner = 160;
124 constexpr uint32_t GM_OUTER = 32;149 constexpr uint32_t GM_OUTER = 32;
125 constexpr uint32_t GM_INNER = 160;150 constexpr uint32_t GM_INNER = 160;
126- op.Init(srcGm, dstGm, lastRowGm, outer, inner, GM_OUTER, GM_INNER, &tPipe);151+ op.Init(srcGm, dstGm, lastRowGm, outer, inner, GM_OUTER, GM_INNER, tmpBufSize, &tPipe);
127 op.Process();152 op.Process();
128}153}
129 154 
@@ -170,6 +195,11 @@ int32_t main(int32_t argc, char* argv[])
170 size_t param4FileSize = 2 * sizeof(uint32_t);195 size_t param4FileSize = 2 * sizeof(uint32_t);
171 uint32_t numBlocks = 1;196 uint32_t numBlocks = 1;
172 197 
198+ ge::Shape shape{{32, 160}};
199+ uint32_t maxValue = 0;
200+ uint32_t minValue = 0;
201+ AscendC::GetCumSumMaxMinTmpSize(shape, sizeof(float), /*isReuseSource*/true, /*isLastAxis*/false, maxValue, minValue);
202+ 
173 aclInit(nullptr);203 aclInit(nullptr);
174 aclrtContext context;204 aclrtContext context;
175 int32_t deviceId = 0;205 int32_t deviceId = 0;
@@ -202,7 +232,7 @@ int32_t main(int32_t argc, char* argv[])
202 ReadFile("./input/input_tiling.bin", param4FileSize, param4Host, param4FileSize);232 ReadFile("./input/input_tiling.bin", param4FileSize, param4Host, param4FileSize);
203 aclrtMemcpy(param4Device, param4FileSize, param4Host, param4FileSize, ACL_MEMCPY_HOST_TO_DEVICE);233 aclrtMemcpy(param4Device, param4FileSize, param4Host, param4FileSize, ACL_MEMCPY_HOST_TO_DEVICE);
204 234 
205- cumsum_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device, param4Device);235+ cumsum_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device, param4Device, minValue);
206 aclrtSynchronizeStream(stream);236 aclrtSynchronizeStream(stream);
207 237 
208 aclrtFree(param1Device);238 aclrtFree(param1Device);
@@ -213,12 +243,10 @@ int32_t main(int32_t argc, char* argv[])
213 aclrtMemcpy(param2Host, param2FileSize, param2Device, param2FileSize, ACL_MEMCPY_DEVICE_TO_HOST);243 aclrtMemcpy(param2Host, param2FileSize, param2Device, param2FileSize, ACL_MEMCPY_DEVICE_TO_HOST);
214 WriteFile("./output/output_result.bin", param2Host, param2FileSize);244 WriteFile("./output/output_result.bin", param2Host, param2FileSize);
215 aclrtFree(param2Device);245 aclrtFree(param2Device);
216- aclrtFreeHost(param2Host);
217 246 
218 aclrtMemcpy(param3Host, param3FileSize, param3Device, param3FileSize, ACL_MEMCPY_DEVICE_TO_HOST);247 aclrtMemcpy(param3Host, param3FileSize, param3Device, param3FileSize, ACL_MEMCPY_DEVICE_TO_HOST);
219 WriteFile("./output/output_last_row.bin", param3Host, param3FileSize);248 WriteFile("./output/output_last_row.bin", param3Host, param3FileSize);
220 aclrtFree(param3Device);249 aclrtFree(param3Device);
221- aclrtFreeHost(param3Host);
222 250 
223 bool goldenResult = true;251 bool goldenResult = true;
224 goldenResult &= CompareResult(param2Host, param2FileSize, "result");252 goldenResult &= CompareResult(param2Host, param2FileSize, "result");
@@ -229,10 +257,13 @@ int32_t main(int32_t argc, char* argv[])
229 printf("test failed!\n");257 printf("test failed!\n");
230 }258 }
231 259 
260+ aclrtFreeHost(param2Host);
261+ aclrtFreeHost(param3Host);
262+ 
232 aclrtDestroyStream(stream);263 aclrtDestroyStream(stream);
233 aclrtDestroyContext(context);264 aclrtDestroyContext(context);
234 aclrtResetDevice(deviceId);265 aclrtResetDevice(deviceId);
235 aclFinalize();266 aclFinalize();
236 267 
237 return 0;268 return 0;
238-}269+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/exp/CMakeLists.txt+6-9
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -25,15 +28,9 @@ target_link_libraries(demo PRIVATE
25 platform28 platform
26 m29 m
27 dl30 dl
31+ graph_base
28)32)
29 33 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE34target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>35+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38- # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>36+)
39-)
Mexamples/01_simd_cpp_api/03_libraries/12_math/exp/README.md+56-25
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Exp高阶API的算子实现。样例按元素取自然指数，用户可以选择是否使用泰勒展开公式进行计算。 5+本样例基于Exp高阶API实现自然指数计算功能，支持按元素计算$e^x$。Exp高阶API可以设置泰勒展开项数，取值范围[0，255]，当泰勒展开项数为0时表示不使用泰勒公式进行计算。本样例默认设置泰勒展开项数为10项。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -12,18 +12,18 @@
12 12 
13## 目录结构介绍13## 目录结构介绍
14 14 
15-```15+```plain
16├── exp16├── exp
17│ ├── scripts17│ ├── scripts
18-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本18+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
19│ ├── CMakeLists.txt // 编译工程文件19│ ├── CMakeLists.txt // 编译工程文件
20│ ├── data_utils.h // 数据读入写出函数20│ ├── data_utils.h // 数据读入写出函数
21│ └── exp.asc // Ascend C算子实现 & 调用样例21│ └── exp.asc // Ascend C算子实现 & 调用样例
22```22```
23 23 
24-## 算子描述24+## 样例描述
25 25 
26-- 算子功能： 26+- 样例功能：
27 按元素取自然指数，用户可以选择是否使用泰勒展开公式进行计算，计算公式如下：27 按元素取自然指数，用户可以选择是否使用泰勒展开公式进行计算，计算公式如下：
28 28 
29 $$29 $$
@@ -32,71 +32,102 @@
32 32 
33 设置泰勒展开项数为0，即不使用泰勒展开公式进行计算，公式如下：33 设置泰勒展开项数为0，即不使用泰勒展开公式进行计算，公式如下：
34 34 
35- $$Exp(x) = e^x$$ 35+ $$Exp(x) = e^x$$
36 36 
37 设置泰勒展开项数不为0，即使用泰勒展开公式进行计算，公式如下：37 设置泰勒展开项数不为0，即使用泰勒展开公式进行计算，公式如下：
38 38 
39 $$Exp(x) = e^{x A_{i}} * e^{x B_{i}}$$39 $$Exp(x) = e^{x A_{i}} * e^{x B_{i}}$$
40 40 
41- $xA_i$代表源操作数的整数部分，该值通过floor(x)获取。xBi代表源操作数的小数部分。41+ $xA_i$代表源操作数的整数部分，该值通过floor(x)获取。$xB_i$代表源操作数的小数部分。
42 42 
43 泰勒展开公式如下：43 泰勒展开公式如下：
44 44 
45 $$e^{x B_i} = 1 + xB_i +... + \frac{xB_i^n}{n!}$$45 $$e^{x B_i} = 1 + xB_i +... + \frac{xB_i^n}{n!}$$
46-
47 46 
48-- 算子规格： 47+- 样例规格：
49 <table>48 <table>
50- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> exp </td></tr>49+ <caption>表1：样例输入输出规格</caption>
50+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> exp </td></tr>
51 51 
52- <tr><td rowspan="3" align="center">算子输入</td></tr>52+ <tr><td rowspan="3" align="center">样例输入</td></tr>
53 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>53 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
54- <tr><td align="center">src</td><td align="center">8192</td><td align="center">float</td><td align="center">ND</td></tr>54+ <tr><td align="center">src</td><td align="center">[1, 8192]</td><td align="center">float</td><td align="center">ND</td></tr>
55- <tr><td rowspan="2" align="center">算子输出</td></tr>55+ <tr><td rowspan="2" align="center">样例输出</td></tr>
56- <tr><td align="center">dst</td><td align="center">8192</td><td align="center">float</td><td align="center">ND</td></tr>56+ <tr><td align="center">dst</td><td align="center">[1, 8192]</td><td align="center">float</td><td align="center">ND</td></tr>
57 57 
58 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">exp_custom</td></tr>58 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">exp_custom</td></tr>
59 </table>59 </table>
60 60 
61-- 算子实现： 61+- 样例实现：
62- 本样例中实现的是固定shape为输入src[8192]，输出dst[8192]的exp_custom算子。62+ 本样例中实现的是固定shape为输入src[1, 8192]，输出dst[1, 8192]的exp_custom样例。
63 63 
64- - Kernel实现 64+ - Kernel实现
65- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Exp高阶API接口完成Exp计算，得到最终结果，再搬出到外部存储上。
66 65 
67- exp_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，Compute任务负责对srcLocal执行Exp计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。66+ 使用Exp高阶API计算自然指数，可选择使用泰勒展开公式和临时buffer提高精度
67+ 
68+ - Tiling实现
69+ 
70+ Host侧通过GetExpMaxMinTmpSize获取Exp接口计算所需的最大和最小临时空间。
71+ 
72+ - 调用实现
68 73 
69- - 调用实现
70 使用内核调用符<<<>>>调用核函数。74 使用内核调用符<<<>>>调用核函数。
71 75 
72## 编译运行 76## 编译运行
73 77 
74在本样例根目录下执行如下步骤，编译并执行算子。78在本样例根目录下执行如下步骤，编译并执行算子。
79+ 
75- 配置环境变量 80- 配置环境变量
76 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。81 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
77 - 默认路径，root用户安装CANN软件包82 - 默认路径，root用户安装CANN软件包
83+ 
78 ```bash84 ```bash
79 source /usr/local/Ascend/cann/set_env.sh85 source /usr/local/Ascend/cann/set_env.sh
80 ```86 ```
81 87 
82 - 默认路径，非root用户安装CANN软件包88 - 默认路径，非root用户安装CANN软件包
89+ 
83 ```bash90 ```bash
84 source $HOME/Ascend/cann/set_env.sh91 source $HOME/Ascend/cann/set_env.sh
85 ```92 ```
86 93 
87 - 指定路径install_path，安装CANN软件包94 - 指定路径install_path，安装CANN软件包
95+ 
88 ```bash96 ```bash
89 source ${install_path}/cann/set_env.sh97 source ${install_path}/cann/set_env.sh
90 ```98 ```
91- 99+ 
92- 样例执行100- 样例执行
101+ 
93 ```bash102 ```bash
94- mkdir -p build && cd build; # 创建并进入build目录103+ mkdir -p build && cd build; # 创建并进入build目录
95- cmake ..;make -j; # 编译工程104+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式
96 python3 ../scripts/gen_data.py # 生成测试输入数据105 python3 ../scripts/gen_data.py # 生成测试输入数据
97- ./demo # 执行编译生成的可执行程序，执行样例106+ ./demo # 执行编译生成的可执行程序，执行样例
98 ```107 ```
108+ 
109+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
110+
111+ 示例如下：
112+ 
113+ ```bash
114+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式
115+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式
116+ ```
117+ 
118+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
119+ 
120+- 编译选项说明
121+ 
122+ | 选项 | 可选值 | 说明 |
123+ |------|--------|------|
124+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
125+ | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2/A3 系列，dav-3510 对应 Ascend 950PR/Ascend 950DT |
126+ 
127+- 执行结果
128+ 
99 执行结果如下，说明精度对比成功。129 执行结果如下，说明精度对比成功。
130+ 
100 ```bash131 ```bash
101 test pass!132 test pass!
102- ```133+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/exp/exp.asc+38-80
@@ -11,86 +11,54 @@
11 11 
12/* !12/* !
13 * \file exp.asc13 * \file exp.asc
14- * \brief14+ * \brief 本样例基于Exp高阶API实现自然指数计算功能，支持按元素计算e^x，可选择使用泰勒展开公式提高精度
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20+#include "tiling/tiling_api.h"
21+ 
22+#ifdef ASCENDC_CPU_DEBUG
23+#include "cpu_debug_launch.h"
24+#endif
20 25 
21constexpr uint32_t SIZE_OF_FLOAT = 4;26constexpr uint32_t SIZE_OF_FLOAT = 4;
22constexpr uint32_t EXP_ONE_BLK_SIZE = 32;27constexpr uint32_t EXP_ONE_BLK_SIZE = 32;
23constexpr uint32_t EXP_MAX_REPEAT = 255;28constexpr uint32_t EXP_MAX_REPEAT = 255;
24constexpr uint32_t EXP_ONE_REPEAT_BYTE_SIZE = 256;29constexpr uint32_t EXP_ONE_REPEAT_BYTE_SIZE = 256;
25 30 
26-__aicore__ inline uint32_t GetFinalBlockSizeTest(uint32_t dataAmount, uint32_t blockSize)
27-{
28- return (dataAmount + blockSize - 1) / blockSize * blockSize;
29-}
30- 
31-__aicore__ inline uint32_t GetExpMaxTmpSizeCur(const uint32_t inputSize, const uint32_t typeSize,
32- const bool isReuseSource, const bool isHighPreci)
33-{
34- uint32_t tmpBufferSize = inputSize * SIZE_OF_FLOAT;
35- if (EXP_ONE_REPEAT_BYTE_SIZE > tmpBufferSize) {
36- tmpBufferSize = EXP_ONE_REPEAT_BYTE_SIZE;
37- }
38- 
39- uint32_t numOfTmpBuf = 4;
40- if (typeSize == 4) {
41- numOfTmpBuf = isReuseSource ? 2 : 3;
42- }
43- return numOfTmpBuf * tmpBufferSize;
44-}
45- 
46-__aicore__ inline uint32_t GetExpMinTmpSizeCur(const uint32_t typeSize, const bool isReuseSource,
47- const bool isHighPreci)
48-{
49- uint32_t numOfTmpBuf = 4;
50- if (typeSize == 4) { // FP32
51- numOfTmpBuf = isReuseSource ? 2 : 3;
52- }
53- return numOfTmpBuf * EXP_ONE_REPEAT_BYTE_SIZE;
54-}
55- 
56-__aicore__ inline bool GetExpMaxMinTmpSizeCur(const AscendC::ShapeInfo& inputShapeInfo, const uint32_t typeSize,
57- const bool isReuseSource, const bool isHighPreci, uint32_t& maxValue,
58- uint32_t& minValue)
59-{
60- uint32_t inputSize = 1;
61- for (int i = 0; i < 1; i++) { inputSize *= inputShapeInfo.shape[i]; }
62- minValue = GetExpMinTmpSizeCur(typeSize, isReuseSource, isHighPreci);
63- maxValue = GetExpMaxTmpSizeCur(inputSize, typeSize, isReuseSource, isHighPreci);
64- return true;
65-}
66 31 
32+/**
33+ * @brief Exp核函数实现类，演示Exp API的使用场景
34+ * @tparam T 数据类型
35+ * @tparam isReuseSrc 是否复用源操作数
36+ * @tparam isHighPreci 是否使用高精度计算
37+ * @tparam expandLevel 泰勒展开项数，为0表示不使用泰勒展开
38+ */
67template <typename T, bool isReuseSrc = false, bool isHighPreci = false, uint8_t expandLevel = 10>39template <typename T, bool isReuseSrc = false, bool isHighPreci = false, uint8_t expandLevel = 10>
68class KernelExp {40class KernelExp {
69public:41public:
70 __aicore__ inline KernelExp() {}42 __aicore__ inline KernelExp() {}
71- __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t totalLength, uint32_t calCount, uint32_t tmpMode,43+ __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t totalLength, uint32_t calCount, uint32_t tmpBufSize, AscendC::TPipe* pipeIn)
72- AscendC::DataFormat dataFormat, AscendC::TPipe* pipeIn)
73 {44 {
74 pipe = pipeIn;45 pipe = pipeIn;
75 this->totalLength = totalLength;46 this->totalLength = totalLength;
76 this->calCount = calCount;47 this->calCount = calCount;
77- this->tmpMode = tmpMode;
78- this->dataFormat = dataFormat;
79 uint32_t oneBlockNum = 32 / sizeof(T);48 uint32_t oneBlockNum = 32 / sizeof(T);
80 totalLength = (totalLength + oneBlockNum - 1) / oneBlockNum * oneBlockNum;49 totalLength = (totalLength + oneBlockNum - 1) / oneBlockNum * oneBlockNum;
81 srcGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(srcGm), totalLength);50 srcGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(srcGm), totalLength);
82 dstGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(dstGm), totalLength);51 dstGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(dstGm), totalLength);
83 pipe->InitBuffer(inQueueX, 1, totalLength * sizeof(T));52 pipe->InitBuffer(inQueueX, 1, totalLength * sizeof(T));
84 pipe->InitBuffer(outQueue, 1, totalLength * sizeof(T));53 pipe->InitBuffer(outQueue, 1, totalLength * sizeof(T));
54+ pipe->InitBuffer(buf, tmpBufSize * sizeof(T));
85 }55 }
86 56 
87 __aicore__ inline void Process()57 __aicore__ inline void Process()
88 {58 {
89- AscendC::AscendCUtils::SetOverflow(1);
90 CopyIn();59 CopyIn();
91 Compute();60 Compute();
92 CopyOut();61 CopyOut();
93- AscendC::AscendCUtils::SetOverflow(0);
94 }62 }
95 63 
96 __aicore__ inline void CopyIn()64 __aicore__ inline void CopyIn()
@@ -104,32 +72,20 @@ public:
104 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();72 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();
105 AscendC::LocalTensor<T> srcLocal = inQueueX.DeQue<T>();73 AscendC::LocalTensor<T> srcLocal = inQueueX.DeQue<T>();
106 Duplicate(dstLocal, (T)0, totalLength);74 Duplicate(dstLocal, (T)0, totalLength);
75+ AscendC::LocalTensor<uint8_t> stackBuffer = buf.Get<uint8_t>();
107 76 
108- uint32_t inputShape[1] = {totalLength};77+ // 使用Exp接口计算自然指数，使用临时buffer提高精度
109- AscendC::ShapeInfo shapeInfo{1, inputShape, 1, inputShape, AscendC::DataFormat::ND};78+ // 模板参数：
110- uint32_t maxSize = 0;79+ // - T: 输入输出数据类型
111- uint32_t minSize = 0;80+ // - expandLevel: 泰勒展开项数，为0表示不使用泰勒展开
112- GetExpMaxMinTmpSizeCur(shapeInfo, sizeof(T), isReuseSrc, isHighPreci, maxSize, minSize);81+ // - isReuseSrc: 是否复用源操作数
82+ // 参数说明：
83+ // - dstLocal: 输出Tensor，存储计算结果
84+ // - srcLocal: 输入Tensor，待计算指数的值
85+ // - stackBuffer: 临时buffer，用于泰勒展开计算
86+ // - calCount: 计算元素个数
87+ AscendC::Exp<T, expandLevel, isReuseSrc>(dstLocal, srcLocal, stackBuffer, calCount);
113 88 
114- AscendC::LocalTensor<uint8_t> stackBuffer;
115- bool ans = AscendC::PopStackBuffer<uint8_t, AscendC::TPosition::LCM>(stackBuffer);
116- stackBufferSize = stackBuffer.GetSize() * 1;
117- 
118- for (int i = 0; i < 1; i++) {
119- if (tmpMode == 0) {
120- AscendC::Exp<T, expandLevel, isReuseSrc>(dstLocal, srcLocal, calCount);
121- } else {
122- if (tmpMode == 1) {
123- stackBufferSize = minSize;
124- } else if (tmpMode == 2) {
125- stackBufferSize = maxSize;
126- } else if (tmpMode == 3) {
127- stackBufferSize = (minSize + maxSize) / 2;
128- }
129- stackBuffer.SetSize(stackBufferSize);
130- AscendC::Exp<T, expandLevel, isReuseSrc>(dstLocal, srcLocal, stackBuffer, calCount);
131- }
132- }
133 outQueue.EnQue<T>(dstLocal);89 outQueue.EnQue<T>(dstLocal);
134 inQueueX.FreeTensor(srcLocal);90 inQueueX.FreeTensor(srcLocal);
135 }91 }
@@ -144,23 +100,20 @@ private:
144 AscendC::TPipe* pipe;100 AscendC::TPipe* pipe;
145 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX;101 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX;
146 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueue;102 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueue;
103+ AscendC::TBuf<AscendC::QuePosition::VECCALC> buf;
147 AscendC::GlobalTensor<T> srcGlobal;104 AscendC::GlobalTensor<T> srcGlobal;
148 AscendC::GlobalTensor<T> dstGlobal;105 AscendC::GlobalTensor<T> dstGlobal;
149 uint32_t totalLength = 0;106 uint32_t totalLength = 0;
150 uint32_t calCount = 0;107 uint32_t calCount = 0;
151- uint32_t tmpMode = 0;
152- AscendC::DataFormat dataFormat;
153- uint32_t stackBufferSize = 0;
154};108};
155 109 
156-__global__ __vector__ void exp_custom(GM_ADDR srcGm, GM_ADDR dstGm)110+__global__ __vector__ void exp_custom(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t tmpBufSize)
157{111{
158 AscendC::TPipe pipe;112 AscendC::TPipe pipe;
159 constexpr uint32_t totalLength = 8192;113 constexpr uint32_t totalLength = 8192;
160 constexpr uint32_t calCount = 8192;114 constexpr uint32_t calCount = 8192;
161- constexpr uint32_t tmpMode = 0;115+ KernelExp<float, false, false> op;
162- KernelExp<float, false, false, 0> op;116+ op.Init(srcGm, dstGm, totalLength, calCount, tmpBufSize, &pipe);
163- op.Init(srcGm, dstGm, totalLength, calCount, tmpMode, AscendC::DataFormat::ND, &pipe);
164 op.Process();117 op.Process();
165}118}
166 119 
@@ -204,6 +157,11 @@ int32_t main(int32_t argc, char* argv[])
204 size_t param2FileSize = 8192 * sizeof(float);157 size_t param2FileSize = 8192 * sizeof(float);
205 uint32_t numBlocks = 1;158 uint32_t numBlocks = 1;
206 159 
160+ ge::Shape shape{{8192}};
161+ uint32_t maxValue = 0;
162+ uint32_t minValue = 0;
163+ AscendC::GetExpMaxMinTmpSize(shape, sizeof(float), false, maxValue, minValue);
164+ 
207 aclInit(nullptr);165 aclInit(nullptr);
208 aclrtContext context;166 aclrtContext context;
209 int32_t deviceId = 0;167 int32_t deviceId = 0;
@@ -224,7 +182,7 @@ int32_t main(int32_t argc, char* argv[])
224 aclrtMallocHost((void**)(&param2Host), param2FileSize);182 aclrtMallocHost((void**)(&param2Host), param2FileSize);
225 aclrtMalloc((void**)&param2Device, param2FileSize, ACL_MEM_MALLOC_HUGE_FIRST);183 aclrtMalloc((void**)&param2Device, param2FileSize, ACL_MEM_MALLOC_HUGE_FIRST);
226 184 
227- exp_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device);185+ exp_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, minValue);
228 aclrtSynchronizeStream(stream);186 aclrtSynchronizeStream(stream);
229 187 
230 aclrtFree(param1Device);188 aclrtFree(param1Device);
@@ -250,4 +208,4 @@ int32_t main(int32_t argc, char* argv[])
250 aclFinalize();208 aclFinalize();
251 209 
252 return 0;210 return 0;
253-}211+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/exp/scripts/gen_data.py+19-10
@@ -16,24 +16,33 @@ import os
16import sys16import sys
17import numpy as np17import numpy as np
18 18 
19+def taylor_exp(src, n):
20+ if n < 1:
21+ raise
22+ item = np.ones_like(src).astype(np.float32)
23+ result = np.ones_like(src).astype(np.float32)
24+ for i in range(n):
25+ item *= src/(i+1)
26+ result += item
27+ return result
19 28 
20-def gen_golden_data_simple():29+def gen_golden_data_simple(item_num=10):
21 dtype = np.float3230 dtype = np.float32
22- src_shape = [8192]31+ cal_count = 8192
32+ src_shape = [cal_count]
23 np.random.seed(0)33 np.random.seed(0)
24 input_dtype = dtype34 input_dtype = dtype
25- cal_count = 8192
26- is_high_preci = 0
27 min_num, max_num = input_dtype(-10), input_dtype(10)35 min_num, max_num = input_dtype(-10), input_dtype(10)
28- element_num = 136+ 
29- for a in src_shape:
30- element_num *= a
31 src = np.random.uniform(min_num, max_num, src_shape).astype(input_dtype)37 src = np.random.uniform(min_num, max_num, src_shape).astype(input_dtype)
32- if is_high_preci == 1:
33- src = src.astype(np.float32)
34 src_exp = src[:cal_count]38 src_exp = src[:cal_count]
35 src_ori = np.zeros(src.size - cal_count).astype(src.dtype)39 src_ori = np.zeros(src.size - cal_count).astype(src.dtype)
36- src_exp = np.exp(src_exp)40+ if item_num:
41+ xa = np.floor(src_exp)
42+ xb = src_exp - xa
43+ src_exp = np.exp(xa) * taylor_exp(xb, item_num)
44+ else:
45+ src_exp = np.exp(src_exp)
37 golden = np.concatenate((src_exp, src_ori), axis=None)46 golden = np.concatenate((src_exp, src_ori), axis=None)
38 47 
39 if input_dtype != np.float32:48 if input_dtype != np.float32:
Mexamples/01_simd_cpp_api/03_libraries/12_math/fma/CMakeLists.txt+4-7
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 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-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -27,12 +30,6 @@ target_link_libraries(demo PRIVATE
27 dl30 dl
28)31)
29 32 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE33target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>34+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38)35)
Mexamples/01_simd_cpp_api/03_libraries/12_math/fma/README.md+56-25
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Fma高阶API的算子实现。样例按元素计算两个输入相乘后与第三个输入相加的结果。5+本样例基于Fma高阶API实现按元素计算两个输入相乘后与第三个输入相加的功能。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -10,82 +10,113 @@
10 10 
11## 目录结构介绍11## 目录结构介绍
12 12 
13-```13+```plain
14├── fma14├── fma
15│ ├── scripts15│ ├── scripts
16-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本16+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
17│ ├── CMakeLists.txt // 编译工程文件17│ ├── CMakeLists.txt // 编译工程文件
18│ ├── data_utils.h // 数据读入写出函数18│ ├── data_utils.h // 数据读入写出函数
19-│ └── fma.asc // Ascend C算子实现 & 调用样例19+│ └── fma.asc // Ascend C样例实现 & 调用样例
20```20```
21 21 
22-## 算子描述22+## 样例描述
23 23 
24-- 算子功能： 24+- 样例功能：
25 按元素计算两个输入相乘后与第三个输入相加的结果。25 按元素计算两个输入相乘后与第三个输入相加的结果。
26 26 
27 计算公式如下：27 计算公式如下：
28 $$28 $$
29 dst_i = Fma(src0_i, src1_i, src2_i)29 dst_i = Fma(src0_i, src1_i, src2_i)
30 $$30 $$
31+ 
31 $$32 $$
32 Fma(src0_i, src1_i, src2_i) = src0_i * src1_i + src2_i33 Fma(src0_i, src1_i, src2_i) = src0_i * src1_i + src2_i
33 $$34 $$
34 35 
35-- 算子规格： 36+- 样例规格：
36 <table>37 <table>
37- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> fma </td></tr>38+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> fma </td></tr>
38 39 
39- <tr><td rowspan="5" align="center">算子输入</td></tr>40+ <tr><td rowspan="5" align="center">样例输入</td></tr>
40 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>41 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
41- <tr><td align="center">src0</td><td align="center">128</td><td align="center">float</td><td align="center">ND</td></tr>42+ <tr><td align="center">src0</td><td align="center">[1, 128]</td><td align="center">float</td><td align="center">ND</td></tr>
42- <tr><td align="center">src1</td><td align="center">128</td><td align="center">float</td><td align="center">ND</td></tr>43+ <tr><td align="center">src1</td><td align="center">[1, 128]</td><td align="center">float</td><td align="center">ND</td></tr>
43- <tr><td align="center">src2</td><td align="center">128</td><td align="center">float</td><td align="center">ND</td></tr>44+ <tr><td align="center">src2</td><td align="center">[1, 128]</td><td align="center">float</td><td align="center">ND</td></tr>
44- <tr><td rowspan="2" align="center">算子输出</td></tr>45+ <tr><td rowspan="2" align="center">样例输出</td></tr>
45- <tr><td align="center">dst</td><td align="center">128</td><td align="center">float</td><td align="center">ND</td></tr>46+ <tr><td align="center">dst</td><td align="center">[1, 128]</td><td align="center">float</td><td align="center">ND</td></tr>
46 47 
47 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">fma_custom</td></tr>48 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">fma_custom</td></tr>
48 </table>49 </table>
49 50 
50-- 算子实现： 51+- 样例实现：
51- 本样例中实现的是固定shape为输入src0[128]、src1[128]、src2[128]，输出dst[128]的fma_custom算子。52+ 本样例中实现的是固定shape为输入src0[1, 128]、src1[1, 128]、src2[1, 128]，输出dst[1, 128]的fma_custom样例。
52 53 
53- - Kernel实现 54+ - Kernel实现
54- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Fma高阶API接口完成Fma计算，得到最终结果，再搬出到外部存储上。
55 55 
56- fma_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor src0Gm、src1Gm、src0Gm存储在srcLocal中，Compute任务负责对src0Local、src1Local、src2Local执行Fma计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。56+ 使用Fma高阶API计算 src0 * src1 + src2，可选择使用临时buffer
57+ 
58+ - Tiling实现
59+ 
60+ Host侧通过GetFmaMaxMinTmpSize获取Fma接口计算所需的最大和最小临时空间。
57 61 
58 - 调用实现 62 - 调用实现
59 使用内核调用符<<<>>>调用核函数。63 使用内核调用符<<<>>>调用核函数。
60 64 
61## 编译运行 65## 编译运行
62 66 
63-在本样例根目录下执行如下步骤，编译并执行算子。67+在本样例根目录下执行如下步骤，编译并执行样例。
68+ 
64- 配置环境变量 69- 配置环境变量
65 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。70 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
66 - 默认路径，root用户安装CANN软件包71 - 默认路径，root用户安装CANN软件包
72+ 
67 ```bash73 ```bash
68 source /usr/local/Ascend/cann/set_env.sh74 source /usr/local/Ascend/cann/set_env.sh
69 ```75 ```
70 76 
71 - 默认路径，非root用户安装CANN软件包77 - 默认路径，非root用户安装CANN软件包
78+ 
72 ```bash79 ```bash
73 source $HOME/Ascend/cann/set_env.sh80 source $HOME/Ascend/cann/set_env.sh
74 ```81 ```
75 82 
76 - 指定路径install_path，安装CANN软件包83 - 指定路径install_path，安装CANN软件包
84+ 
77 ```bash85 ```bash
78 source ${install_path}/cann/set_env.sh86 source ${install_path}/cann/set_env.sh
79 ```87 ```
80- 88+ 
81- 样例执行89- 样例执行
90+ 
82 ```bash91 ```bash
83- mkdir -p build && cd build; # 创建并进入build目录92+ mkdir -p build && cd build; # 创建并进入build目录
84- cmake ..;make -j; # 编译工程93+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程，默认npu模式
85 python3 ../scripts/gen_data.py # 生成测试输入数据94 python3 ../scripts/gen_data.py # 生成测试输入数据
86- ./demo # 执行编译生成的可执行程序，执行样例95+ ./demo # 执行编译生成的可执行程序，执行样例
87 ```96 ```
97+ 
98+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
99+
100+ 示例如下：
101+ 
102+ ```bash
103+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
104+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
105+ ```
106+ 
107+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
108+ 
109+- 编译选项说明
110+ 
111+ | 选项 | 可选值 | 说明 |
112+ |------|--------|------|
113+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
114+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认） | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
115+ 
116+- 执行结果
117+ 
88 执行结果如下，说明精度对比成功。118 执行结果如下，说明精度对比成功。
119+ 
89 ```bash120 ```bash
90 test pass!121 test pass!
91- ```122+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/fma/fma.asc+38-14
@@ -11,18 +11,29 @@
11 11 
12/* !12/* !
13 * \file fma.asc13 * \file fma.asc
14- * \brief14+ * \brief 本样例基于Fma高阶API实现乘加融合运算功能，按元素计算两个输入相乘后与第三个输入相加的结果
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20 20 
21-template <typename T, int32_t calCount, int32_t dataSize, int32_t sharedTmpBufferSize>21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
25+/**
26+ * @brief Fma核函数实现类，演示Fma API的使用场景
27+ * @tparam T 数据类型
28+ * @tparam calCount 计算元素个数
29+ * @tparam dataSize 数据大小
30+ * @tparam sharedTmpBufferSize 临时buffer大小
31+ */
32+template <typename T, int32_t calCount, int32_t dataSize>
22class KernelFma {33class KernelFma {
23public:34public:
24 __aicore__ inline KernelFma() {}35 __aicore__ inline KernelFma() {}
25- __aicore__ inline void Init(GM_ADDR src0Gm, GM_ADDR src1Gm, GM_ADDR src2Gm, GM_ADDR dstGm, AscendC::TPipe* pipeIn)36+ __aicore__ inline void Init(GM_ADDR src0Gm, GM_ADDR src1Gm, GM_ADDR src2Gm, GM_ADDR dstGm, uint32_t tmpBufSize, AscendC::TPipe* pipeIn)
26 {37 {
27 pipe = pipeIn;38 pipe = pipeIn;
28 src0Global.SetGlobalBuffer((__gm__ T*)(src0Gm));39 src0Global.SetGlobalBuffer((__gm__ T*)(src0Gm));
@@ -34,17 +45,15 @@ public:
34 pipe->InitBuffer(inQueue1, 1, dataSize * sizeof(T));45 pipe->InitBuffer(inQueue1, 1, dataSize * sizeof(T));
35 pipe->InitBuffer(inQueue2, 1, dataSize * sizeof(T));46 pipe->InitBuffer(inQueue2, 1, dataSize * sizeof(T));
36 pipe->InitBuffer(outQueue, 1, dataSize * sizeof(T));47 pipe->InitBuffer(outQueue, 1, dataSize * sizeof(T));
37- if constexpr (sharedTmpBufferSize > 0) {48+ if (tmpBufSize > 0) {
38- pipe->InitBuffer(bufQueue, sharedTmpBufferSize * sizeof(T));49+ pipe->InitBuffer(bufQueue, tmpBufSize * sizeof(T));
39 }50 }
40 }51 }
41 __aicore__ inline void Process()52 __aicore__ inline void Process()
42 {53 {
43- AscendC::AscendCUtils::SetOverflow(1);
44 CopyIn();54 CopyIn();
45 Compute();55 Compute();
46 CopyOut();56 CopyOut();
47- AscendC::AscendCUtils::SetOverflow(0);
48 }57 }
49 58 
50 __aicore__ inline void CopyIn()59 __aicore__ inline void CopyIn()
@@ -68,7 +77,16 @@ public:
68 AscendC::LocalTensor<T> src1Local = inQueue1.DeQue<T>();77 AscendC::LocalTensor<T> src1Local = inQueue1.DeQue<T>();
69 AscendC::LocalTensor<T> src2Local = inQueue2.DeQue<T>();78 AscendC::LocalTensor<T> src2Local = inQueue2.DeQue<T>();
70 AscendC::Duplicate(dstLocal, (T)0, dataSize);79 AscendC::Duplicate(dstLocal, (T)0, dataSize);
71- if constexpr (sharedTmpBufferSize > 0) {80+ // 使用Fma接口计算乘加融合运算
81+ // 模板参数：
82+ // - T: 输入输出数据类型
83+ // 参数说明：
84+ // - dstLocal: 输出Tensor，存储计算结果
85+ // - src0Local: 第一个输入Tensor
86+ // - src1Local: 第二个输入Tensor
87+ // - src2Local: 第三个输入Tensor
88+ // - calCount: 计算元素个数
89+ if (tmpBufSize > 0) {
72 AscendC::LocalTensor<uint8_t> tmpBuf = bufQueue.Get<uint8_t>();90 AscendC::LocalTensor<uint8_t> tmpBuf = bufQueue.Get<uint8_t>();
73 AscendC::Fma(dstLocal, src0Local, src1Local, src2Local, tmpBuf, calCount);91 AscendC::Fma(dstLocal, src0Local, src1Local, src2Local, tmpBuf, calCount);
74 } else {92 } else {
@@ -99,14 +117,13 @@ private:
99 AscendC::GlobalTensor<T> dstGlobal;117 AscendC::GlobalTensor<T> dstGlobal;
100};118};
101 119 
102-__global__ __vector__ void fma_custom(GM_ADDR src0Gm, GM_ADDR src1Gm, GM_ADDR src2Gm, GM_ADDR dstGm)120+__global__ __vector__ void fma_custom(GM_ADDR src0Gm, GM_ADDR src1Gm, GM_ADDR src2Gm, GM_ADDR dstGm, uint32_t tmpBufSize)
103{121{
104 AscendC::TPipe pipe;122 AscendC::TPipe pipe;
105 constexpr uint32_t dataSize = 128;123 constexpr uint32_t dataSize = 128;
106 constexpr uint32_t calCount = 128;124 constexpr uint32_t calCount = 128;
107- constexpr uint32_t sharedTmpBufSize = 0;125+ KernelFma<float, calCount, dataSize> op;
108- KernelFma<float, calCount, dataSize, sharedTmpBufSize> op;126+ op.Init(src0Gm, src1Gm, src2Gm, dstGm, tmpBufSize, &pipe);
109- op.Init(src0Gm, src1Gm, src2Gm, dstGm, &pipe);
110 op.Process();127 op.Process();
111}128}
112 129 
@@ -152,6 +169,13 @@ int32_t main(int32_t argc, char* argv[])
152 size_t param4FileSize = 128 * sizeof(float);169 size_t param4FileSize = 128 * sizeof(float);
153 uint32_t numBlocks = 1;170 uint32_t numBlocks = 1;
154 171 
172+ platform_ascendc::PlatformAscendC* ascendcPlatform = platform_ascendc::PlatformAscendCManager::GetInstance();
173+ const platform_ascendc::PlatformAscendC& plat = *ascendcPlatform;
174+ ge::Shape shape{{128}};
175+ uint32_t maxValue = 0;
176+ uint32_t minValue = 0;
177+ AscendC::GetRintMaxMinTmpSize(plat, shape, sizeof(float), false, maxValue, minValue);
178+ 
155 aclInit(nullptr);179 aclInit(nullptr);
156 aclrtContext context;180 aclrtContext context;
157 int32_t deviceId = 0;181 int32_t deviceId = 0;
@@ -186,7 +210,7 @@ int32_t main(int32_t argc, char* argv[])
186 aclrtMallocHost((void**)(&param4Host), param4FileSize);210 aclrtMallocHost((void**)(&param4Host), param4FileSize);
187 aclrtMalloc((void**)&param4Device, param4FileSize, ACL_MEM_MALLOC_HUGE_FIRST);211 aclrtMalloc((void**)&param4Device, param4FileSize, ACL_MEM_MALLOC_HUGE_FIRST);
188 212 
189- fma_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device, param4Device);213+ fma_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device, param4Device, minValue);
190 aclrtSynchronizeStream(stream);214 aclrtSynchronizeStream(stream);
191 215 
192 aclrtFree(param1Device);216 aclrtFree(param1Device);
@@ -216,4 +240,4 @@ int32_t main(int32_t argc, char* argv[])
216 aclFinalize();240 aclFinalize();
217 241 
218 return 0;242 return 0;
219-}243+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/fmod/CMakeLists.txt+4-8
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -27,13 +30,6 @@ target_link_libraries(demo PRIVATE
27 dl30 dl
28)31)
29 32 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE33target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>34+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38- # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>
39)35)
Mexamples/01_simd_cpp_api/03_libraries/12_math/fmod/README.md+54-24
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Fmod高阶API的算子实现。样例按元素计算两个浮点数a，b相除后的余数。5+本样例基于Fmod高阶API实现按元素浮点数取余的功能。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -12,18 +12,18 @@
12 12 
13## 目录结构介绍13## 目录结构介绍
14 14 
15-```15+```plain
16├── fmod16├── fmod
17│ ├── scripts17│ ├── scripts
18-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本18+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
19│ ├── CMakeLists.txt // 编译工程文件19│ ├── CMakeLists.txt // 编译工程文件
20│ ├── data_utils.h // 数据读入写出函数20│ ├── data_utils.h // 数据读入写出函数
21-│ └── fmod.asc // Ascend C算子实现 & 调用样例21+│ └── fmod.asc // Ascend C样例实现 & 调用样例
22```22```
23 23 
24-## 算子描述24+## 样例描述
25 25 
26-- 算子功能： 26+- 样例功能：
27 按元素计算两个浮点数a，b相除后的余数。27 按元素计算两个浮点数a，b相除后的余数。
28 28 
29 计算公式如下：29 计算公式如下：
@@ -37,59 +37,89 @@
37 37 
38 Fmod(-3.0, 1.1) = -0.838 Fmod(-3.0, 1.1) = -0.8
39 39 
40-- 算子规格： 40+- 样例规格：
41 <table>41 <table>
42- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> fmod </td></tr>42+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> fmod </td></tr>
43 43 
44- <tr><td rowspan="4" align="center">算子输入</td></tr>44+ <tr><td rowspan="4" align="center">样例输入</td></tr>
45 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>45 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
46- <tr><td align="center">src0</td><td align="center">159</td><td align="center">float</td><td align="center">ND</td></tr>46+ <tr><td align="center">src0</td><td align="center">[1, 159]</td><td align="center">float</td><td align="center">ND</td></tr>
47- <tr><td align="center">src1</td><td align="center">159</td><td align="center">float</td><td align="center">ND</td></tr>47+ <tr><td align="center">src1</td><td align="center">[1, 159]</td><td align="center">float</td><td align="center">ND</td></tr>
48- <tr><td rowspan="2" align="center">算子输出</td></tr>48+ <tr><td rowspan="2" align="center">样例输出</td></tr>
49- <tr><td align="center">dst</td><td align="center">159</td><td align="center">float</td><td align="center">ND</td></tr>49+ <tr><td align="center">dst</td><td align="center">[1, 159]</td><td align="center">float</td><td align="center">ND</td></tr>
50 50 
51 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">fmod_custom</td></tr>51 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">fmod_custom</td></tr>
52 </table>52 </table>
53 53 
54-- 算子实现： 54+- 样例实现：
55- 本样例中实现的是固定shape为输入src0[159]、src1[159]，输出dst[159]的fmod_custom算子。55+ 本样例中实现的是固定shape为输入src0[1, 159]、src1[1, 159]，输出dst[1, 159]的fmod_custom样例。
56 56 
57- - Kernel实现 57+ - Kernel实现
58- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Fmod高阶API接口完成Fmod计算，得到最终结果，再搬出到外部存储上。
59 58 
60- fmod_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor src0Gm、src1Gm存储在srcLocal中，Compute任务负责对src0Local、src1Local执行Fmod计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。59+ 使用Fmod高阶API计算取余运算，可选择使用临时buffer和指定计算元素个数
60+ 
61+ - Tiling实现
62+ 
63+ Host侧通过GetFmodMaxMinTmpSize获取Fmod接口计算所需的最大和最小临时空间。
61 64 
62 - 调用实现 65 - 调用实现
63 使用内核调用符<<<>>>调用核函数。66 使用内核调用符<<<>>>调用核函数。
64 67 
65## 编译运行 68## 编译运行
66 69 
67-在本样例根目录下执行如下步骤，编译并执行算子。70+在本样例根目录下执行如下步骤，编译并执行样例。
71+ 
68- 配置环境变量 72- 配置环境变量
69 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。73 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
70 - 默认路径，root用户安装CANN软件包74 - 默认路径，root用户安装CANN软件包
75+ 
71 ```bash76 ```bash
72 source /usr/local/Ascend/cann/set_env.sh77 source /usr/local/Ascend/cann/set_env.sh
73 ```78 ```
74 79 
75 - 默认路径，非root用户安装CANN软件包80 - 默认路径，非root用户安装CANN软件包
81+ 
76 ```bash82 ```bash
77 source $HOME/Ascend/cann/set_env.sh83 source $HOME/Ascend/cann/set_env.sh
78 ```84 ```
79 85 
80 - 指定路径install_path，安装CANN软件包86 - 指定路径install_path，安装CANN软件包
87+ 
81 ```bash88 ```bash
82 source ${install_path}/cann/set_env.sh89 source ${install_path}/cann/set_env.sh
83 ```90 ```
84- 91+ 
85- 样例执行92- 样例执行
93+ 
86 ```bash94 ```bash
87- mkdir -p build && cd build; # 创建并进入build目录95+ mkdir -p build && cd build; # 创建并进入build目录
88- cmake ..;make -j; # 编译工程96+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式
89 python3 ../scripts/gen_data.py # 生成测试输入数据97 python3 ../scripts/gen_data.py # 生成测试输入数据
90- ./demo # 执行编译生成的可执行程序，执行样例98+ ./demo # 执行编译生成的可执行程序，执行样例
91 ```99 ```
100+ 
101+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
102+
103+ 示例如下：
104+ 
105+ ```bash
106+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式
107+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式
108+ ```
109+ 
110+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
111+ 
112+- 编译选项说明
113+ 
114+ | 选项 | 可选值 | 说明 |
115+ |------|--------|------|
116+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
117+ | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2/A3 系列，dav-3510 对应 Ascend 950PR/Ascend 950DT |
118+ 
119+- 执行结果
120+ 
92 执行结果如下，说明精度对比成功。121 执行结果如下，说明精度对比成功。
122+ 
93 ```bash123 ```bash
94 test pass!124 test pass!
95- ```125+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/fmod/fmod.asc+32-6
@@ -11,13 +11,17 @@
11 11 
12/* !12/* !
13 * \file fmod.asc13 * \file fmod.asc
14- * \brief14+ * \brief 本样例基于Fmod高阶API实现取余运算功能，按元素计算两个浮点数相除后的余数
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20 20 
21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
21constexpr int32_t BUFFER_NUM = 1;25constexpr int32_t BUFFER_NUM = 1;
22 26 
23template <typename T>27template <typename T>
@@ -27,6 +31,13 @@ __aicore__ inline uint32_t Align32B(uint32_t len)
27 return (len + alginSize - 1) / alginSize * alginSize;31 return (len + alginSize - 1) / alginSize * alginSize;
28}32}
29 33 
34+/**
35+ * @brief Fmod核函数实现类，演示Fmod API的使用场景
36+ * @tparam T 数据类型
37+ * @tparam IS_REUSE_SOURCE 是否复用源操作数
38+ * @tparam USE_SHARED_TMP_BUFFER 是否使用临时buffer
39+ * @tparam USE_CAL_COUNT 是否使用计算元素个数参数
40+ */
30template <typename T, bool IS_REUSE_SOURCE, bool USE_SHARED_TMP_BUFFER, bool USE_CAL_COUNT>41template <typename T, bool IS_REUSE_SOURCE, bool USE_SHARED_TMP_BUFFER, bool USE_CAL_COUNT>
31class KernelFmod {42class KernelFmod {
32public:43public:
@@ -76,6 +87,17 @@ public:
76 AscendC::LocalTensor<T> src0Local = src0Queue.DeQue<T>();87 AscendC::LocalTensor<T> src0Local = src0Queue.DeQue<T>();
77 AscendC::LocalTensor<T> src1Local = src1Queue.DeQue<T>();88 AscendC::LocalTensor<T> src1Local = src1Queue.DeQue<T>();
78 89 
90+ // 使用Fmod接口计算取余运算
91+ // 模板参数：
92+ // - T: 输入输出数据类型
93+ // - IS_REUSE_SOURCE: 是否复用源操作数
94+ // - config: Fmod配置参数（仅3510架构）
95+ // 参数说明：
96+ // - dstLocal: 输出Tensor，存储计算结果
97+ // - src0Local: 被除数Tensor
98+ // - src1Local: 除数Tensor
99+ // - sharedTmpBuffer: 临时buffer，用于提高精度
100+ // - calCount: 计算元素个数
79#if __NPU_ARCH__ == 3510101#if __NPU_ARCH__ == 3510
80 static constexpr AscendC::FmodConfig config = {AscendC::FmodAlgo::NORMAL, AscendC::FMOD_ITERATION_NUM_MAX};102 static constexpr AscendC::FmodConfig config = {AscendC::FmodAlgo::NORMAL, AscendC::FMOD_ITERATION_NUM_MAX};
81 if constexpr (USE_SHARED_TMP_BUFFER) {103 if constexpr (USE_SHARED_TMP_BUFFER) {
@@ -137,15 +159,14 @@ private:
137 uint32_t sharedTmpBufferSize{1};159 uint32_t sharedTmpBufferSize{1};
138};160};
139 161 
140-__vector__ __global__ void fmod_custom(GM_ADDR src0Gm, GM_ADDR src1Gm, GM_ADDR dstGm)162+__vector__ __global__ void fmod_custom(GM_ADDR src0Gm, GM_ADDR src1Gm, GM_ADDR dstGm, uint32_t tmpBufSize)
141{163{
142 AscendC::TPipe pipe;164 AscendC::TPipe pipe;
143 constexpr uint32_t inCount = 159;165 constexpr uint32_t inCount = 159;
144 constexpr uint32_t outCount = 159;166 constexpr uint32_t outCount = 159;
145 constexpr uint32_t calCount = 159;167 constexpr uint32_t calCount = 159;
146- constexpr uint32_t bufferSize = 2000;
147 KernelFmod<float, 0, 0, 1> op;168 KernelFmod<float, 0, 0, 1> op;
148- op.Init(src0Gm, src1Gm, dstGm, inCount, outCount, calCount, bufferSize, &pipe);169+ op.Init(src0Gm, src1Gm, dstGm, inCount, outCount, calCount, tmpBufSize, &pipe);
149 op.Process();170 op.Process();
150}171}
151 172 
@@ -190,6 +211,11 @@ int32_t main(int32_t argc, char* argv[])
190 size_t param3FileSize = 159 * sizeof(float);211 size_t param3FileSize = 159 * sizeof(float);
191 uint32_t numBlocks = 1;212 uint32_t numBlocks = 1;
192 213 
214+ ge::Shape shape{{159}};
215+ uint32_t maxValue = 0;
216+ uint32_t minValue = 0;
217+ AscendC::GetFmodMaxMinTmpSize(shape, sizeof(float), false, maxValue, minValue);
218+ 
193 aclInit(nullptr);219 aclInit(nullptr);
194 aclrtContext context;220 aclrtContext context;
195 int32_t deviceId = 0;221 int32_t deviceId = 0;
@@ -217,7 +243,7 @@ int32_t main(int32_t argc, char* argv[])
217 aclrtMallocHost((void**)(&param3Host), param3FileSize);243 aclrtMallocHost((void**)(&param3Host), param3FileSize);
218 aclrtMalloc((void**)&param3Device, param3FileSize, ACL_MEM_MALLOC_HUGE_FIRST);244 aclrtMalloc((void**)&param3Device, param3FileSize, ACL_MEM_MALLOC_HUGE_FIRST);
219 245 
220- fmod_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device);246+ fmod_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device, minValue);
221 aclrtSynchronizeStream(stream);247 aclrtSynchronizeStream(stream);
222 248 
223 aclrtFree(param1Device);249 aclrtFree(param1Device);
@@ -245,4 +271,4 @@ int32_t main(int32_t argc, char* argv[])
245 aclFinalize();271 aclFinalize();
246 272 
247 return 0;273 return 0;
248-}274+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/frac/CMakeLists.txt+6-9
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -25,15 +28,9 @@ target_link_libraries(demo PRIVATE
25 platform28 platform
26 m29 m
27 dl30 dl
31+ graph_base
28)32)
29 33 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE34target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>35+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38- # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>36+)
39-)
Mexamples/01_simd_cpp_api/03_libraries/12_math/frac/README.md+54-24
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Frac高阶API的算子实现。样例按元素做取小数计算。5+本样例基于Frac高阶API实现按元素取小数的功能。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -12,73 +12,103 @@
12 12 
13## 目录结构介绍13## 目录结构介绍
14 14 
15-```15+```plain
16├── frac16├── frac
17│ ├── scripts17│ ├── scripts
18-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本18+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
19│ ├── CMakeLists.txt // 编译工程文件19│ ├── CMakeLists.txt // 编译工程文件
20│ ├── data_utils.h // 数据读入写出函数20│ ├── data_utils.h // 数据读入写出函数
21-│ └── frac.asc // Ascend C算子实现 & 调用样例21+│ └── frac.asc // Ascend C样例实现 & 调用样例
22```22```
23 23 
24-## 算子描述24+## 样例描述
25 25 
26-- 算子功能： 26+- 样例功能：
27- 按元素做双曲正弦函数计算，计算公式如下： 27+ 按元素取小数，计算公式如下：
28 $$dstTensor_i = Frac(srcTensor_i)$$28 $$dstTensor_i = Frac(srcTensor_i)$$
29 29 
30-- 算子规格： 30+- 样例规格：
31 <table>31 <table>
32- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> frac </td></tr>32+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> frac </td></tr>
33 33 
34- <tr><td rowspan="3" align="center">算子输入</td></tr>34+ <tr><td rowspan="3" align="center">样例输入</td></tr>
35 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>35 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
36- <tr><td align="center">src</td><td align="center">4096</td><td align="center">float</td><td align="center">ND</td></tr>36+ <tr><td align="center">src</td><td align="center">[1, 4096]</td><td align="center">float</td><td align="center">ND</td></tr>
37- <tr><td rowspan="2" align="center">算子输出</td></tr>37+ <tr><td rowspan="2" align="center">样例输出</td></tr>
38- <tr><td align="center">dst</td><td align="center">4096</td><td align="center">float</td><td align="center">ND</td></tr>38+ <tr><td align="center">dst</td><td align="center">[1, 4096]</td><td align="center">float</td><td align="center">ND</td></tr>
39 39 
40 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">frac_custom</td></tr>40 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">frac_custom</td></tr>
41 </table>41 </table>
42 42 
43-- 算子实现： 43+- 样例实现：
44- 本样例中实现的是固定shape为输入src[4096]，输出dst[4096]的frac_custom算子。44+ 本样例中实现的是固定shape为输入src[1, 4096]，输出dst[1, 4096]的frac_custom样例。
45 45 
46- - Kernel实现 46+ - Kernel实现
47- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Frac高阶API接口完成Frac计算，得到最终结果，再搬出到外部存储上。
48 47 
49- frac_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，Compute任务负责对srcLocal执行Frac计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。48+ 使用Frac高阶API按元素取小数部分
49+ 
50+ - Tiling实现
51+ 
52+ Host侧通过GetFracMaxMinTmpSize获取Frac接口计算所需的最大和最小临时空间。
50 53 
51 - 调用实现 54 - 调用实现
52 使用内核调用符<<<>>>调用核函数。55 使用内核调用符<<<>>>调用核函数。
53 56 
54## 编译运行 57## 编译运行
55 58 
56-在本样例根目录下执行如下步骤，编译并执行算子。59+在本样例根目录下执行如下步骤，编译并执行样例。
60+ 
57- 配置环境变量 61- 配置环境变量
58 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。62 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
59 - 默认路径，root用户安装CANN软件包63 - 默认路径，root用户安装CANN软件包
64+ 
60 ```bash65 ```bash
61 source /usr/local/Ascend/cann/set_env.sh66 source /usr/local/Ascend/cann/set_env.sh
62 ```67 ```
63 68 
64 - 默认路径，非root用户安装CANN软件包69 - 默认路径，非root用户安装CANN软件包
70+ 
65 ```bash71 ```bash
66 source $HOME/Ascend/cann/set_env.sh72 source $HOME/Ascend/cann/set_env.sh
67 ```73 ```
68 74 
69 - 指定路径install_path，安装CANN软件包75 - 指定路径install_path，安装CANN软件包
76+ 
70 ```bash77 ```bash
71 source ${install_path}/cann/set_env.sh78 source ${install_path}/cann/set_env.sh
72 ```79 ```
73- 80+ 
74- 样例执行81- 样例执行
82+ 
75 ```bash83 ```bash
76- mkdir -p build && cd build; # 创建并进入build目录84+ mkdir -p build && cd build; # 创建并进入build目录
77- cmake ..;make -j; # 编译工程85+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式
78 python3 ../scripts/gen_data.py # 生成测试输入数据86 python3 ../scripts/gen_data.py # 生成测试输入数据
79- ./demo # 执行编译生成的可执行程序，执行样例87+ ./demo # 执行编译生成的可执行程序，执行样例
80 ```88 ```
89+ 
90+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
91+
92+ 示例如下：
93+ 
94+ ```bash
95+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式
96+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式
97+ ```
98+ 
99+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
100+ 
101+- 编译选项说明
102+ 
103+ | 选项 | 可选值 | 说明 |
104+ |------|--------|------|
105+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
106+ | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2/A3 系列，dav-3510 对应 Ascend 950PR/Ascend 950DT |
107+ 
108+- 执行结果
109+ 
81 执行结果如下，说明精度对比成功。110 执行结果如下，说明精度对比成功。
111+ 
82 ```bash112 ```bash
83 test pass!113 test pass!
84- ```114+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/frac/frac.asc+33-14
@@ -11,18 +11,28 @@
11 11 
12/* !12/* !
13 * \file frac.asc13 * \file frac.asc
14- * \brief14+ * \brief 本样例基于Frac高阶API实现取小数部分功能，按元素提取浮点数的小数部分
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20+#include "tiling/tiling_api.h"
20 21 
22+#ifdef ASCENDC_CPU_DEBUG
23+#include "cpu_debug_launch.h"
24+#endif
25+ 
26+/**
27+ * @brief Frac核函数实现类，演示Frac API的使用场景
28+ * @tparam T 数据类型
29+ * @tparam apiMode API调用模式：0-无参数，1-带calCount，2-带tmpBuf和calCount，3-带tmpBuf
30+ */
21template <typename T, int32_t apiMode>31template <typename T, int32_t apiMode>
22class KernelFrac {32class KernelFrac {
23public:33public:
24 __aicore__ inline KernelFrac() {}34 __aicore__ inline KernelFrac() {}
25- __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t srcSize, uint32_t calCount,35+ __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t srcSize, uint32_t calCount, uint32_t tmpBufSize,
26 AscendC::TPipe* pipeIn)36 AscendC::TPipe* pipeIn)
27 {37 {
28 pipe = pipeIn;38 pipe = pipeIn;
@@ -35,34 +45,39 @@ public:
35 alignDataNum = alignDataSize / sizeof(T);45 alignDataNum = alignDataSize / sizeof(T);
36 46 
37 pipe->InitBuffer(inQueueX, 1, alignDataSize);47 pipe->InitBuffer(inQueueX, 1, alignDataSize);
38- pipe->InitBuffer(inQueueT, 1, alignDataNum * sizeof(uint8_t));
39 pipe->InitBuffer(outQueue, 1, alignDataSize);48 pipe->InitBuffer(outQueue, 1, alignDataSize);
49+ pipe->InitBuffer(buf, tmpBufSize * sizeof(uint8_t));
40 }50 }
41 __aicore__ inline void Process()51 __aicore__ inline void Process()
42 {52 {
43- AscendC::AscendCUtils::SetOverflow(1);
44 CopyIn();53 CopyIn();
45 Compute();54 Compute();
46 CopyOut();55 CopyOut();
47- AscendC::AscendCUtils::SetOverflow(0);
48 }56 }
49 57 
50 __aicore__ inline void CopyIn()58 __aicore__ inline void CopyIn()
51 {59 {
52 AscendC::LocalTensor<T> srcLocal = inQueueX.AllocTensor<T>();60 AscendC::LocalTensor<T> srcLocal = inQueueX.AllocTensor<T>();
53- AscendC::LocalTensor<uint8_t> tmpBufLocal = inQueueT.AllocTensor<uint8_t>();
54 AscendC::DataCopy(srcLocal, srcGlobal, bufferSize);61 AscendC::DataCopy(srcLocal, srcGlobal, bufferSize);
55 inQueueX.EnQue(srcLocal);62 inQueueX.EnQue(srcLocal);
56- inQueueT.EnQue(tmpBufLocal);
57 }63 }
58 __aicore__ inline void Compute()64 __aicore__ inline void Compute()
59 {65 {
60 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();66 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();
61 AscendC::LocalTensor<T> srcLocal = inQueueX.DeQue<T>();67 AscendC::LocalTensor<T> srcLocal = inQueueX.DeQue<T>();
62- AscendC::LocalTensor<uint8_t> tmpBuf = inQueueT.DeQue<uint8_t>();68+ AscendC::LocalTensor<uint8_t> tmpBuf = buf.Get<uint8_t>();
63 T zero = 0;69 T zero = 0;
64 AscendC::Duplicate(dstLocal, zero, alignDataNum);70 AscendC::Duplicate(dstLocal, zero, alignDataNum);
65 71 
72+ // 使用Frac接口提取小数部分
73+ // 模板参数：
74+ // - T: 输入输出数据类型
75+ // - false: 是否复用源操作数
76+ // 参数说明：
77+ // - dstLocal: 输出Tensor，存储计算结果
78+ // - srcLocal: 输入Tensor
79+ // - tmpBuf: 临时buffer，用于提高精度
80+ // - count: 计算元素个数
66 if constexpr (apiMode == 0) {81 if constexpr (apiMode == 0) {
67 AscendC::Frac<T, false>(dstLocal, srcLocal);82 AscendC::Frac<T, false>(dstLocal, srcLocal);
D
DDragonBornHD844月25日

更重要的模板参数没做介绍，入参稍微看看代码还能看得懂，模板参数不做注释说明是真不知道什么含义

likedislike
68 } else if constexpr (apiMode == 1) {83 } else if constexpr (apiMode == 1) {
@@ -75,7 +90,6 @@ public:
75 90 
76 outQueue.EnQue<T>(dstLocal);91 outQueue.EnQue<T>(dstLocal);
77 inQueueX.FreeTensor(srcLocal);92 inQueueX.FreeTensor(srcLocal);
78- inQueueT.FreeTensor(tmpBuf);
79 }93 }
80 __aicore__ inline void CopyOut()94 __aicore__ inline void CopyOut()
81 {95 {
@@ -87,8 +101,8 @@ public:
87private:101private:
88 AscendC::TPipe* pipe;102 AscendC::TPipe* pipe;
89 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX;103 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX;
90- AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueT;
91 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueue;104 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueue;
105+ AscendC::TBuf<AscendC::QuePosition::VECCALC> buf;
92 AscendC::GlobalTensor<T> srcGlobal;106 AscendC::GlobalTensor<T> srcGlobal;
93 AscendC::GlobalTensor<T> dstGlobal;107 AscendC::GlobalTensor<T> dstGlobal;
94 uint32_t bufferSize = 0;108 uint32_t bufferSize = 0;
@@ -97,14 +111,14 @@ private:
97 uint32_t count = 0;111 uint32_t count = 0;
98};112};
99 113 
100-__global__ __vector__ void frac_custom(GM_ADDR srcGm, GM_ADDR dstGm)114+__global__ __vector__ void frac_custom(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t tmpBufSize)
101{115{
102 AscendC::TPipe pipe;116 AscendC::TPipe pipe;
103 constexpr uint32_t srcSize = 4096;117 constexpr uint32_t srcSize = 4096;
104 constexpr uint32_t calCount = 4096;118 constexpr uint32_t calCount = 4096;
105 constexpr uint32_t apiMode = 0;119 constexpr uint32_t apiMode = 0;
106 KernelFrac<float, apiMode> op;120 KernelFrac<float, apiMode> op;
107- op.Init(srcGm, dstGm, srcSize, calCount, &pipe);121+ op.Init(srcGm, dstGm, srcSize, calCount, tmpBufSize, &pipe);
108 op.Process();122 op.Process();
109}123}
110 124 
@@ -148,6 +162,11 @@ int32_t main(int32_t argc, char* argv[])
148 size_t param2FileSize = 4096 * sizeof(float);162 size_t param2FileSize = 4096 * sizeof(float);
149 uint32_t numBlocks = 1;163 uint32_t numBlocks = 1;
150 164 
165+ ge::Shape shape{{4096}};
166+ uint32_t maxValue = 0;
167+ uint32_t minValue = 0;
168+ AscendC::GetFracMaxMinTmpSize(shape, sizeof(float), false, maxValue, minValue);
169+ 
151 aclInit(nullptr);170 aclInit(nullptr);
152 aclrtContext context;171 aclrtContext context;
153 int32_t deviceId = 0;172 int32_t deviceId = 0;
@@ -168,7 +187,7 @@ int32_t main(int32_t argc, char* argv[])
168 aclrtMallocHost((void**)(&param2Host), param2FileSize);187 aclrtMallocHost((void**)(&param2Host), param2FileSize);
169 aclrtMalloc((void**)&param2Device, param2FileSize, ACL_MEM_MALLOC_HUGE_FIRST);188 aclrtMalloc((void**)&param2Device, param2FileSize, ACL_MEM_MALLOC_HUGE_FIRST);
170 189 
171- frac_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device);190+ frac_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, minValue);
172 aclrtSynchronizeStream(stream);191 aclrtSynchronizeStream(stream);
173 192 
174 aclrtFree(param1Device);193 aclrtFree(param1Device);
@@ -193,4 +212,4 @@ int32_t main(int32_t argc, char* argv[])
193 aclFinalize();212 aclFinalize();
194 213 
195 return 0;214 return 0;
196-}215+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/power/CMakeLists.txt+5-9
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -27,13 +30,6 @@ target_link_libraries(demo PRIVATE
27 dl30 dl
28)31)
29 32 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE33target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>34+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38- # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>35+)
39-)
Mexamples/01_simd_cpp_api/03_libraries/12_math/power/README.md+60-27
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Power高阶API的算子实现。样例实现按元素做幂运算功能，支持三种功能：指数和底数分别为张量对张量、张量对标量、标量对张量的幂运算。5+本样例基于Power高阶API实现按元素做幂运算功能，支持三种功能：指数和底数分别为张量对张量、张量对标量、标量对张量的幂运算。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -12,82 +12,115 @@
12 12 
13## 目录结构介绍13## 目录结构介绍
14 14 
15-```15+```plain
16├── power16├── power
17│ ├── scripts17│ ├── scripts
18-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本18+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
19│ ├── CMakeLists.txt // 编译工程文件19│ ├── CMakeLists.txt // 编译工程文件
20│ ├── data_utils.h // 数据读入写出函数20│ ├── data_utils.h // 数据读入写出函数
21-│ └── power.asc // Ascend C算子实现 & 调用样例21+│ └── power.asc // Ascend C样例实现 & 调用样例
22```22```
23 23 
24-## 算子描述24+## 样例描述
25 25 
26-- 算子功能： 26+- 样例功能：
27- 实现按元素做幂运算功能，支持三种功能：指数和底数分别为张量对张量、张量对标量、标量对张量的幂运算，参数mode分别为0，1，2。27+ 实现按元素做幂运算功能，支持三种功能：指数和底数分别为张量对张量、张量对标量、标量对张量的幂运算。
28 28 
29 计算公式如下：29 计算公式如下：
30 $$Power(x, y) = x^y$$30 $$Power(x, y) = x^y$$
31- 张量对张量，mode = 0：两个长度相同的张量，逐元素做幂运算 31+ 
32+ 张量对张量，mode = 0：两个长度相同的张量，逐元素做幂运算
32 $$dstTensor_i = Power(srcbaseTensor_i, srcexpTensor_i)$$33 $$dstTensor_i = Power(srcbaseTensor_i, srcexpTensor_i)$$
34+ 
33 张量对标量， mode = 1：以标量作为指数，张量都用同一个指数进行幂运算35 张量对标量， mode = 1：以标量作为指数，张量都用同一个指数进行幂运算
34 $$dstTensor_i = Power(srcbaseTensor_i, srcexpScalar)$$36 $$dstTensor_i = Power(srcbaseTensor_i, srcexpScalar)$$
35- 标量对张量， mode = 2：以标量作为固定的底数，张量都用同一个底数进行幂运算 37+ 
38+ 标量对张量， mode = 2：以标量作为固定的底数，张量都用同一个底数进行幂运算
36 $$dstTensor_i = Power(srcbaseScalar, srcexpTensor_i)$$39 $$dstTensor_i = Power(srcbaseScalar, srcexpTensor_i)$$
37 40 
38-- 算子规格： 41+- 样例规格：
39 <table>42 <table>
40- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> power </td></tr>43+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> power </td></tr>
41 44 
42- <tr><td rowspan="4" align="center">算子输入</td></tr>45+ <tr><td rowspan="4" align="center">样例输入</td></tr>
43 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>46 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
44- <tr><td align="center">srcbase</td><td align="center">16</td><td align="center">float</td><td align="center">ND</td></tr>47+ <tr><td align="center">srcbase</td><td align="center">[1, 16]</td><td align="center">float</td><td align="center">ND</td></tr>
45- <tr><td align="center">srcexp</td><td align="center">16</td><td align="center">float</td><td align="center">ND</td></tr>48+ <tr><td align="center">srcexp</td><td align="center">[1, 16]</td><td align="center">float</td><td align="center">ND</td></tr>
46- <tr><td rowspan="2" align="center">算子输出</td></tr>49+ <tr><td rowspan="2" align="center">样例输出</td></tr>
47- <tr><td align="center">dst</td><td align="center">16</td><td align="center">float</td><td align="center">ND</td></tr>50+ <tr><td align="center">dst</td><td align="center">[1, 16]</td><td align="center">float</td><td align="center">ND</td></tr>
48 51 
49 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">power_custom</td></tr>52 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">power_custom</td></tr>
50 </table>53 </table>
51 54 
52-- 算子实现： 55+- 样例实现：
53- 本样例中实现的是固定shape为输入srcbase[16]、srcexp[16]，输出dst[16]的power_custom算子。算子功能mode参数默认为0，即指数和底数都为张量。56+ 本样例中实现的是固定shape为输入srcbase[1, 16]、srcexp[1, 16]，输出dst[1, 16]的power_custom样例。样例功能mode参数默认为0，即指数和底数都为张量。
54 57 
55- - Kernel实现 58+ - Kernel实现
56- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Power高阶API接口完成Power计算，得到最终结果，再搬出到外部存储上。
57 59 
58- power_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcbaseGm、srcexpGm存储在srcLocal中，输入参数mode用于判断幂运算的类型，Compute任务负责对srcbaseLocal、srcexpLocal执行Power计算，按照mode的值调用不同类调用接口实现计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。60+ 使用Power高阶API进行幂运算，支持张量对张量、张量对标量、标量对张量三种模式
61+ 
62+ - Tiling实现
63+ 
64+ Host侧通过GetPowerMaxMinTmpSize获取Power接口计算所需的最大和最小临时空间。
59 65 
60 - 调用实现 66 - 调用实现
61 使用内核调用符<<<>>>调用核函数。67 使用内核调用符<<<>>>调用核函数。
62 68 
63## 编译运行 69## 编译运行
64 70 
65-在本样例根目录下执行如下步骤，编译并执行算子。71+在本样例根目录下执行如下步骤，编译并执行样例。
72+ 
66- 配置环境变量 73- 配置环境变量
67 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。74 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
68 - 默认路径，root用户安装CANN软件包75 - 默认路径，root用户安装CANN软件包
76+ 
69 ```bash77 ```bash
70 source /usr/local/Ascend/cann/set_env.sh78 source /usr/local/Ascend/cann/set_env.sh
71 ```79 ```
72 80 
73 - 默认路径，非root用户安装CANN软件包81 - 默认路径，非root用户安装CANN软件包
82+ 
74 ```bash83 ```bash
75 source $HOME/Ascend/cann/set_env.sh84 source $HOME/Ascend/cann/set_env.sh
76 ```85 ```
77 86 
78 - 指定路径install_path，安装CANN软件包87 - 指定路径install_path，安装CANN软件包
88+ 
79 ```bash89 ```bash
80 source ${install_path}/cann/set_env.sh90 source ${install_path}/cann/set_env.sh
81 ```91 ```
82- 92+ 
83- 样例执行93- 样例执行
94+ 
84 ```bash95 ```bash
85- mkdir -p build && cd build; # 创建并进入build目录96+ mkdir -p build && cd build; # 创建并进入build目录
86- cmake ..;make -j; # 编译工程97+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式
87 python3 ../scripts/gen_data.py # 生成测试输入数据98 python3 ../scripts/gen_data.py # 生成测试输入数据
88- ./demo # 执行编译生成的可执行程序，执行样例99+ ./demo # 执行编译生成的可执行程序，执行样例
89 ```100 ```
101+ 
102+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
103+
104+ 示例如下：
105+ 
106+ ```bash
107+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式
108+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式
109+ ```
110+ 
111+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
112+ 
113+- 编译选项说明
114+ 
115+ | 选项 | 可选值 | 说明 |
116+ |------|--------|------|
117+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
118+ | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2/A3 系列，dav-3510 对应 Ascend 950PR/Ascend 950DT |
119+ 
120+- 执行结果
121+ 
90 执行结果如下，说明精度对比成功。122 执行结果如下，说明精度对比成功。
123+ 
91 ```bash124 ```bash
92 test pass!125 test pass!
93- ```126+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/power/power.asc+50-16
@@ -11,18 +11,26 @@
11 11 
12/* !12/* !
13 * \file power.asc13 * \file power.asc
14- * \brief14+ * \brief 本样例基于Power高阶API实现幂运算功能，支持张量对张量、张量对标量、标量对张量三种模式
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20 20 
21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
25+/**
26+ * @brief Power核函数实现类，演示Power API的使用场景
27+ * @tparam T 数据类型
28+ */
21template <typename T>29template <typename T>
22class KernelPower {30class KernelPower {
23public:31public:
24 __aicore__ inline KernelPower() {}32 __aicore__ inline KernelPower() {}
25- __aicore__ inline void Init(GM_ADDR srcGmBase, GM_ADDR srcGmExp, GM_ADDR dstGm, uint32_t srcSize, uint32_t mode,33+ __aicore__ inline void Init(GM_ADDR srcGmBase, GM_ADDR srcGmExp, GM_ADDR dstGm, uint32_t srcSize, uint32_t tmpBufSize,
26 AscendC::TPipe* pipeIn)34 AscendC::TPipe* pipeIn)
27 {35 {
28 pipe = pipeIn;36 pipe = pipeIn;
@@ -33,16 +41,14 @@ public:
33 pipe->InitBuffer(inQueueX1, 1, srcSize * sizeof(T));41 pipe->InitBuffer(inQueueX1, 1, srcSize * sizeof(T));
34 pipe->InitBuffer(inQueueX2, 1, srcSize * sizeof(T));42 pipe->InitBuffer(inQueueX2, 1, srcSize * sizeof(T));
35 pipe->InitBuffer(outQueue, 1, srcSize * sizeof(T));43 pipe->InitBuffer(outQueue, 1, srcSize * sizeof(T));
44+ pipe->InitBuffer(buf, tmpBufSize);
36 bufferSize = srcSize;45 bufferSize = srcSize;
37- this->mode = mode;
38 }46 }
39 __aicore__ inline void Process()47 __aicore__ inline void Process()
40 {48 {
41- AscendC::AscendCUtils::SetOverflow(1);
42 CopyIn();49 CopyIn();
43 Compute();50 Compute();
44 CopyOut();51 CopyOut();
45- AscendC::AscendCUtils::SetOverflow(0);
46 }52 }
47 53 
48 __aicore__ inline void CopyIn()54 __aicore__ inline void CopyIn()
@@ -59,34 +65,48 @@ public:
59 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();65 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();
60 AscendC::LocalTensor<T> srcLocalBase = inQueueX1.DeQue<T>();66 AscendC::LocalTensor<T> srcLocalBase = inQueueX1.DeQue<T>();
61 AscendC::LocalTensor<T> srcLocalExp = inQueueX2.DeQue<T>();67 AscendC::LocalTensor<T> srcLocalExp = inQueueX2.DeQue<T>();
68+ AscendC::LocalTensor<uint8_t> tmpBuffer = buf.Get<uint8_t>();
62 69 
70+ // 使用Power接口进行幂运算
71+ // 模板参数：
72+ // - T: 输入输出数据类型
73+ // - false: 是否复用源操作数
74+ // - config: Power配置参数（仅3510架构）
75+ // 参数说明：
76+ // - dstLocal: 输出Tensor，存储计算结果
77+ // - srcLocalBase: 底数Tensor
78+ // - srcLocalExp: 指数Tensor
79+ // mode参数说明：
80+ // - 0: 张量对张量，逐元素做幂运算
81+ // - 1: 张量对标量，以标量作为指数
82+ // - 2: 标量对张量，以标量作为底数
63#if __NPU_ARCH__ == 351083#if __NPU_ARCH__ == 3510
64 static constexpr AscendC::PowerAlgo valueLowPrecision = AscendC::PowerAlgo::INTRINSIC;84 static constexpr AscendC::PowerAlgo valueLowPrecision = AscendC::PowerAlgo::INTRINSIC;
65 static constexpr AscendC::PowerAlgo valueHighPrecision = AscendC::PowerAlgo::DOUBLE_FLOAT_TECH;85 static constexpr AscendC::PowerAlgo valueHighPrecision = AscendC::PowerAlgo::DOUBLE_FLOAT_TECH;
66 static constexpr AscendC::PowerConfig configLowPrecision = {valueLowPrecision};86 static constexpr AscendC::PowerConfig configLowPrecision = {valueLowPrecision};
67 static constexpr AscendC::PowerConfig configHighPrecision = {valueHighPrecision};87 static constexpr AscendC::PowerConfig configHighPrecision = {valueHighPrecision};
68 if (mode == 0) {88 if (mode == 0) {
69- AscendC::Power<T, false, configHighPrecision>(dstLocal, srcLocalBase, srcLocalExp);89+ AscendC::Power<T, false, configHighPrecision>(dstLocal, srcLocalBase, srcLocalExp, tmpBuffer);
70 } else if (mode == 1) {90 } else if (mode == 1) {
71 T scalarValue = srcLocalExp.GetValue(0);91 T scalarValue = srcLocalExp.GetValue(0);
72 AscendC::PipeBarrier<PIPE_V>();92 AscendC::PipeBarrier<PIPE_V>();
73- AscendC::Power<T, false, configHighPrecision>(dstLocal, srcLocalBase, scalarValue);93+ AscendC::Power<T, false, configHighPrecision>(dstLocal, srcLocalBase, scalarValue, tmpBuffer);
74 } else if (mode == 2) {94 } else if (mode == 2) {
75 T scalarValue = srcLocalBase.GetValue(0);95 T scalarValue = srcLocalBase.GetValue(0);
76 AscendC::PipeBarrier<PIPE_V>();96 AscendC::PipeBarrier<PIPE_V>();
77- AscendC::Power<T, false, configHighPrecision>(dstLocal, scalarValue, srcLocalExp);97+ AscendC::Power<T, false, configHighPrecision>(dstLocal, scalarValue, srcLocalExp, tmpBuffer);
78 }98 }
79#elif __NPU_ARCH__ == 220199#elif __NPU_ARCH__ == 2201
80 if (mode == 0) {100 if (mode == 0) {
81- AscendC::Power<T, false>(dstLocal, srcLocalBase, srcLocalExp);101+ AscendC::Power<T, false>(dstLocal, srcLocalBase, srcLocalExp, tmpBuffer);
82 } else if (mode == 1) {102 } else if (mode == 1) {
83 T scalarValue = srcLocalExp.GetValue(0);103 T scalarValue = srcLocalExp.GetValue(0);
84 AscendC::PipeBarrier<PIPE_V>();104 AscendC::PipeBarrier<PIPE_V>();
85- AscendC::Power<T, false>(dstLocal, srcLocalBase, scalarValue);105+ AscendC::Power<T, false>(dstLocal, srcLocalBase, scalarValue, tmpBuffer);
86 } else if (mode == 2) {106 } else if (mode == 2) {
87 T scalarValue = srcLocalBase.GetValue(0);107 T scalarValue = srcLocalBase.GetValue(0);
88 AscendC::PipeBarrier<PIPE_V>();108 AscendC::PipeBarrier<PIPE_V>();
89- AscendC::Power<T, false>(dstLocal, scalarValue, srcLocalExp);109+ AscendC::Power<T, false>(dstLocal, scalarValue, srcLocalExp, tmpBuffer);
90 }110 }
91#endif111#endif
92 AscendC::PipeBarrier<PIPE_V>();112 AscendC::PipeBarrier<PIPE_V>();
@@ -110,17 +130,18 @@ private:
110 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX1;130 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX1;
111 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX2;131 AscendC::TQue<AscendC::QuePosition::VECIN, 1> inQueueX2;
112 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueue;132 AscendC::TQue<AscendC::QuePosition::VECOUT, 1> outQueue;
133+ AscendC::TBuf<AscendC::TPosition::VECCALC> buf;
113 134 
114 uint32_t bufferSize = 0;135 uint32_t bufferSize = 0;
115- uint32_t mode = 0;136+ uint32_t mode = SCENARIO;
116};137};
117 138 
118__vector__ __global__ void power_custom(GM_ADDR srcGmBase, GM_ADDR srcGmExp, GM_ADDR dstGm, uint32_t srcSize,139__vector__ __global__ void power_custom(GM_ADDR srcGmBase, GM_ADDR srcGmExp, GM_ADDR dstGm, uint32_t srcSize,
119- uint32_t mode)140+ uint32_t tmpBufSize)
120{141{
121 AscendC::TPipe pipe;142 AscendC::TPipe pipe;
122 KernelPower<float> op;143 KernelPower<float> op;
123- op.Init(srcGmBase, srcGmExp, dstGm, srcSize, mode, &pipe);144+ op.Init(srcGmBase, srcGmExp, dstGm, srcSize, tmpBufSize, &pipe);
124 op.Process();145 op.Process();
125}146}
126 147 
@@ -167,6 +188,19 @@ int32_t main(int32_t argc, char* argv[])
167 constexpr uint32_t mode = 0; // 默认模式为mode=0，即指数和底数都为张量188 constexpr uint32_t mode = 0; // 默认模式为mode=0，即指数和底数都为张量
168 uint32_t numBlocks = 1;189 uint32_t numBlocks = 1;
169 190 
191+ ge::Shape shape1{{16}};
192+ ge::Shape shape2{{16}};
193+ ge::Shape shapeScalar{{1}};
194+ uint32_t maxValue = 0;
195+ uint32_t minValue = 0;
196+ if (SCENARIO == 0){
197+ AscendC::GetPowerMaxMinTmpSize(shape1, shape2, false, sizeof(float), false, maxValue, minValue);
198+ } else if (SCENARIO == 1) {
199+ AscendC::GetPowerMaxMinTmpSize(shape1, shapeScalar, false, sizeof(float), false, maxValue, minValue);
200+ } else if (SCENARIO == 2) {
201+ AscendC::GetPowerMaxMinTmpSize(shapeScalar, shape2, false, sizeof(float), false, maxValue, minValue);
202+ }
203+ 
170 aclInit(nullptr);204 aclInit(nullptr);
171 aclrtContext context;205 aclrtContext context;
172 int32_t deviceId = 0;206 int32_t deviceId = 0;
@@ -194,7 +228,7 @@ int32_t main(int32_t argc, char* argv[])
194 aclrtMallocHost((void**)(&param3Host), param3FileSize);228 aclrtMallocHost((void**)(&param3Host), param3FileSize);
195 aclrtMalloc((void**)&param3Device, param3FileSize, ACL_MEM_MALLOC_HUGE_FIRST);229 aclrtMalloc((void**)&param3Device, param3FileSize, ACL_MEM_MALLOC_HUGE_FIRST);
196 230 
197- power_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device, srcSize, mode);231+ power_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device, srcSize, minValue);
198 aclrtSynchronizeStream(stream);232 aclrtSynchronizeStream(stream);
199 233 
200 aclrtFree(param1Device);234 aclrtFree(param1Device);
@@ -222,4 +256,4 @@ int32_t main(int32_t argc, char* argv[])
222 aclFinalize();256 aclFinalize();
223 257 
224 return 0;258 return 0;
225-}259+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/rint/CMakeLists.txt+5-8
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 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-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -27,12 +30,6 @@ target_link_libraries(demo PRIVATE
27 dl30 dl
28)31)
29 32 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE33target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>34+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38-)35+)
Mexamples/01_simd_cpp_api/03_libraries/12_math/rint/README.md+53-23
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Rint高阶API的算子实现。样例获取与输入数据最接近的整数，若存在两个相同接近的整数，则获取其中的偶数。5+本样例基于Rint高阶API实现获取与输入数据最接近整数的功能，若存在两个相同接近的整数，则获取其中的偶数。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -10,18 +10,18 @@
10 10 
11## 目录结构介绍11## 目录结构介绍
12 12 
13-```13+```plain
14├── rint14├── rint
15│ ├── scripts15│ ├── scripts
16-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本16+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
17│ ├── CMakeLists.txt // 编译工程文件17│ ├── CMakeLists.txt // 编译工程文件
18│ ├── data_utils.h // 数据读入写出函数18│ ├── data_utils.h // 数据读入写出函数
19-│ └── rint.asc // Ascend C算子实现 & 调用样例19+│ └── rint.asc // Ascend C样例实现 & 调用样例
20```20```
21 21 
22-## 算子描述22+## 样例描述
23 23 
24-- 算子功能： 24+- 样例功能：
25 获取与输入数据最接近的整数，若存在两个相同接近的整数，则获取其中的偶数。25 获取与输入数据最接近的整数，若存在两个相同接近的整数，则获取其中的偶数。
26 26 
27 计算公式如下：27 计算公式如下：
@@ -29,58 +29,88 @@
29 dst_i = Rint(src_i)29 dst_i = Rint(src_i)
30 $$30 $$
31 31 
32-- 算子规格： 32+- 样例规格：
33 <table>33 <table>
34- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> rint </td></tr>34+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> rint </td></tr>
35 35 
36- <tr><td rowspan="3" align="center">算子输入</td></tr>36+ <tr><td rowspan="3" align="center">样例输入</td></tr>
37 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>37 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
38- <tr><td align="center">src</td><td align="center">1024</td><td align="center">float</td><td align="center">ND</td></tr>38+ <tr><td align="center">src</td><td align="center">[1, 1024]</td><td align="center">float</td><td align="center">ND</td></tr>
39- <tr><td rowspan="2" align="center">算子输出</td></tr>39+ <tr><td rowspan="2" align="center">样例输出</td></tr>
40- <tr><td align="center">dst</td><td align="center">1024</td><td align="center">float</td><td align="center">ND</td></tr>40+ <tr><td align="center">dst</td><td align="center">[1, 1024]</td><td align="center">float</td><td align="center">ND</td></tr>
41 41 
42 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">rint_custom</td></tr>42 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">rint_custom</td></tr>
43 </table>43 </table>
44 44 
45-- 算子实现： 45+- 样例实现：
D
DDragonBornHD844月25日

为啥rint又不加tmpsize的tiling函数的用法了？

likedislike
46- 本样例中实现的是固定shape为输入src[1024]，输出dst[1024]的rint_custom算子。46+ 本样例中实现的是固定shape为输入src[1, 1024]，输出dst[1, 1024]的rint_custom样例。
47 47 
48- - Kernel实现 48+ - Kernel实现
49- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Rint高阶API接口完成Rint计算，得到最终结果，再搬出到外部存储上。
50 49 
51- rint_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor srcGm存储在srcLocal中，Compute任务负责对srcLocal执行Rint计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。50+ 使用Rint高阶API获取与输入数据最接近的整数，若存在两个相同接近的整数则取偶数
51+ 
52+ - Tiling实现
53+ 
54+ Host侧通过GetRintMaxMinTmpSize获取Rint接口计算所需的最大和最小临时空间。
52 55 
53 - 调用实现 56 - 调用实现
54 使用内核调用符<<<>>>调用核函数。57 使用内核调用符<<<>>>调用核函数。
55 58 
56## 编译运行 59## 编译运行
57 60 
58-在本样例根目录下执行如下步骤，编译并执行算子。61+在本样例根目录下执行如下步骤，编译并执行样例。
62+ 
59- 配置环境变量 63- 配置环境变量
60 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。64 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
61 - 默认路径，root用户安装CANN软件包65 - 默认路径，root用户安装CANN软件包
66+ 
62 ```bash67 ```bash
63 source /usr/local/Ascend/cann/set_env.sh68 source /usr/local/Ascend/cann/set_env.sh
64 ```69 ```
65 70 
66 - 默认路径，非root用户安装CANN软件包71 - 默认路径，非root用户安装CANN软件包
72+ 
67 ```bash73 ```bash
68 source $HOME/Ascend/cann/set_env.sh74 source $HOME/Ascend/cann/set_env.sh
69 ```75 ```
70 76 
71 - 指定路径install_path，安装CANN软件包77 - 指定路径install_path，安装CANN软件包
78+ 
72 ```bash79 ```bash
73 source ${install_path}/cann/set_env.sh80 source ${install_path}/cann/set_env.sh
74 ```81 ```
75- 82+ 
76- 样例执行83- 样例执行
84+ 
77 ```bash85 ```bash
78- mkdir -p build && cd build; # 创建并进入build目录86+ mkdir -p build && cd build; # 创建并进入build目录
79- cmake ..;make -j; # 编译工程87+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程，默认npu模式
80 python3 ../scripts/gen_data.py # 生成测试输入数据88 python3 ../scripts/gen_data.py # 生成测试输入数据
81- ./demo # 执行编译生成的可执行程序，执行样例89+ ./demo # 执行编译生成的可执行程序，执行样例
82 ```90 ```
91+ 
92+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
93+
94+ 示例如下：
95+ 
96+ ```bash
97+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
98+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
99+ ```
100+ 
101+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
102+ 
103+- 编译选项说明
104+ 
105+ | 选项 | 可选值 | 说明 |
106+ |------|--------|------|
107+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
108+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认） | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
109+ 
110+- 执行结果
111+ 
83 执行结果如下，说明精度对比成功。112 执行结果如下，说明精度对比成功。
113+ 
84 ```bash114 ```bash
85 test pass!115 test pass!
86- ```116+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/rint/rint.asc+34-12
@@ -11,35 +11,44 @@
11 11 
12/* !12/* !
13 * \file rint.asc13 * \file rint.asc
14- * \brief14+ * \brief 本样例基于Rint高阶API实现四舍五入到最近整数功能，若存在两个相同接近的整数则取偶数
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20 20 
21-template <typename T, int32_t calCount, int32_t dataSize, int32_t sharedTmpBufferSize>21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
25+/**
26+ * @brief Rint核函数实现类，演示Rint API的使用场景
27+ * @tparam T 数据类型
28+ * @tparam calCount 计算元素个数
29+ * @tparam dataSize 数据大小
30+ * @tparam sharedTmpBufferSize 临时buffer大小
31+ */
32+template <typename T, int32_t calCount, int32_t dataSize>
22class KernelRint {33class KernelRint {
23public:34public:
24 __aicore__ inline KernelRint() {}35 __aicore__ inline KernelRint() {}
25- __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, AscendC::TPipe* pipeIn)36+ __aicore__ inline void Init(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t tmpBufSize, AscendC::TPipe* pipeIn)
26 {37 {
27 pipe = pipeIn;38 pipe = pipeIn;
28 srcGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(srcGm));39 srcGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(srcGm));
29 dstGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(dstGm));40 dstGlobal.SetGlobalBuffer(reinterpret_cast<__gm__ T*>(dstGm));
30 pipe->InitBuffer(inQueue, 1, dataSize * sizeof(T));41 pipe->InitBuffer(inQueue, 1, dataSize * sizeof(T));
31 pipe->InitBuffer(outQueue, 1, dataSize * sizeof(T));42 pipe->InitBuffer(outQueue, 1, dataSize * sizeof(T));
32- if constexpr (sharedTmpBufferSize > 0) {43+ if (tmpBufSize > 0) {
33- pipe->InitBuffer(bufQueue, sharedTmpBufferSize * sizeof(T));44+ pipe->InitBuffer(bufQueue, tmpBufSize * sizeof(T));
34 }45 }
35 }46 }
36 __aicore__ inline void Process()47 __aicore__ inline void Process()
37 {48 {
38- AscendC::AscendCUtils::SetOverflow(1);
39 CopyIn();49 CopyIn();
40 Compute();50 Compute();
41 CopyOut();51 CopyOut();
42- AscendC::AscendCUtils::SetOverflow(0);
43 }52 }
44 53 
45 __aicore__ inline void CopyIn()54 __aicore__ inline void CopyIn()
@@ -53,7 +62,13 @@ public:
53 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();62 AscendC::LocalTensor<T> dstLocal = outQueue.AllocTensor<T>();
54 AscendC::LocalTensor<T> srcLocal = inQueue.DeQue<T>();63 AscendC::LocalTensor<T> srcLocal = inQueue.DeQue<T>();
55 AscendC::Duplicate(dstLocal, (T)0, dataSize);64 AscendC::Duplicate(dstLocal, (T)0, dataSize);
56- if constexpr (sharedTmpBufferSize > 0) {65+ // 使用Rint接口四舍五入到最近的偶数
66+ // 参数说明：
67+ // - dstLocal: 输出Tensor，存储计算结果
68+ // - srcLocal: 输入Tensor
69+ // - tmpBuf: 临时buffer，用于提高精度（可选）
70+ // - calCount: 计算元素个数
71+ if (tmpBufSize > 0) {
57 AscendC::LocalTensor<uint8_t> tmpBuf = bufQueue.Get<uint8_t>();72 AscendC::LocalTensor<uint8_t> tmpBuf = bufQueue.Get<uint8_t>();
58 AscendC::Rint(dstLocal, srcLocal, tmpBuf, calCount);73 AscendC::Rint(dstLocal, srcLocal, tmpBuf, calCount);
59 } else {74 } else {
@@ -78,14 +93,14 @@ private:
78 AscendC::GlobalTensor<T> dstGlobal;93 AscendC::GlobalTensor<T> dstGlobal;
79};94};
80 95 
81-__global__ __vector__ void rint_custom(GM_ADDR srcGm, GM_ADDR dstGm)96+__global__ __vector__ void rint_custom(GM_ADDR srcGm, GM_ADDR dstGm, uint32_t tmpBufSize)
82{97{
83 AscendC::TPipe pipe;98 AscendC::TPipe pipe;
84 constexpr uint32_t dataSize = 1024;99 constexpr uint32_t dataSize = 1024;
85 constexpr uint32_t calCount = 1024;100 constexpr uint32_t calCount = 1024;
86 constexpr uint32_t sharedTmpBufSize = 1024;101 constexpr uint32_t sharedTmpBufSize = 1024;
87 KernelRint<float, calCount, (dataSize + 31) / 32 * 32, sharedTmpBufSize> op;102 KernelRint<float, calCount, (dataSize + 31) / 32 * 32, sharedTmpBufSize> op;
88- op.Init(srcGm, dstGm, &pipe);103+ op.Init(srcGm, dstGm, tmpBufSize, &pipe);
89 op.Process();104 op.Process();
90}105}
91 106 
@@ -129,6 +144,13 @@ int32_t main(int32_t argc, char* argv[])
129 size_t param2FileSize = 1024 * sizeof(int32_t);144 size_t param2FileSize = 1024 * sizeof(int32_t);
130 uint32_t numBlocks = 1;145 uint32_t numBlocks = 1;
131 146 
147+ platform_ascendc::PlatformAscendC* ascendcPlatform = platform_ascendc::PlatformAscendCManager::GetInstance();
148+ const platform_ascendc::PlatformAscendC& plat = *ascendcPlatform;
149+ ge::Shape shape{{16}};
150+ uint32_t maxValue = 0;
151+ uint32_t minValue = 0;
152+ AscendC::GetRintMaxMinTmpSize(plat, shape, sizeof(int32_t), false, maxValue, minValue);
153+ 
132 aclInit(nullptr);154 aclInit(nullptr);
133 aclrtContext context;155 aclrtContext context;
134 int32_t deviceId = 0;156 int32_t deviceId = 0;
@@ -149,7 +171,7 @@ int32_t main(int32_t argc, char* argv[])
149 aclrtMallocHost((void**)(&param2Host), param2FileSize);171 aclrtMallocHost((void**)(&param2Host), param2FileSize);
150 aclrtMalloc((void**)&param2Device, param2FileSize, ACL_MEM_MALLOC_HUGE_FIRST);172 aclrtMalloc((void**)&param2Device, param2FileSize, ACL_MEM_MALLOC_HUGE_FIRST);
151 173 
152- rint_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device);174+ rint_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, minValue);
153 aclrtSynchronizeStream(stream);175 aclrtSynchronizeStream(stream);
154 176 
155 aclrtFree(param1Device);177 aclrtFree(param1Device);
@@ -175,4 +197,4 @@ int32_t main(int32_t argc, char* argv[])
175 aclFinalize();197 aclFinalize();
176 198 
177 return 0;199 return 0;
178-}200+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/where/CMakeLists.txt+4-7
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 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-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -19,12 +22,6 @@ add_executable(demo
19 where.asc22 where.asc
20)23)
21 24 
22-# ======================================================================================
23-# NPU 编译选项配置
24-#
25-# 说明：
26-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
27-# ======================================================================================
28target_compile_options(demo PRIVATE25target_compile_options(demo PRIVATE
29- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>26+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
30)27)
Mexamples/01_simd_cpp_api/03_libraries/12_math/where/README.md+53-27
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Where高阶API的算子实现。样例根据指定的条件，从两个源操作数中选择元素，生成目标操作数。两个源操作数均可以是LocalTensor或标量。 5+本样例基于Where高阶API实现根据指定的条件从两个源操作数中选择元素的功能。两个源操作数均可以是LocalTensor或标量。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -10,80 +10,106 @@
10 10 
11## 目录结构介绍11## 目录结构介绍
12 12 
13-```13+```plain
14├── where14├── where
15│ ├── scripts15│ ├── scripts
16-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本16+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
17│ ├── CMakeLists.txt // 编译工程文件17│ ├── CMakeLists.txt // 编译工程文件
18│ ├── data_utils.h // 数据读入写出函数18│ ├── data_utils.h // 数据读入写出函数
19-│ └── where.asc // Ascend C算子实现 & 调用样例19+│ └── where.asc // Ascend C样例实现 & 调用样例
20```20```
21 21 
22-## 算子描述22+## 样例描述
23 23 
24-- 算子功能： 24+- 样例功能：
25 根据指定的条件，从两个源操作数中选择元素，生成目标操作数。两个源操作数均可以是LocalTensor或标量。25 根据指定的条件，从两个源操作数中选择元素，生成目标操作数。两个源操作数均可以是LocalTensor或标量。
26- 26+ 
27 计算公式如下： 27 计算公式如下：
28 $$dst_i = \begin{cases}28 $$dst_i = \begin{cases}
29 src0, & if condition \\29 src0, & if condition \\
30- src1, & otherwise 30+ src1, & otherwise
31 \end{cases}$$31 \end{cases}$$
32 32 
33-- 算子规格： 33+- 样例规格：
34 <table>34 <table>
35- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> where </td></tr>35+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> where </td></tr>
36 36 
37- <tr><td rowspan="5" align="center">算子输入</td></tr>37+ <tr><td rowspan="5" align="center">样例输入</td></tr>
38 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>38 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
39- <tr><td align="center">src0</td><td align="center">32</td><td align="center">float</td><td align="center">ND</td></tr>39+ <tr><td align="center">src0</td><td align="center">[1, 32]</td><td align="center">float</td><td align="center">ND</td></tr>
40- <tr><td align="center">src1</td><td align="center">32</td><td align="center">float</td><td align="center">ND</td></tr>40+ <tr><td align="center">src1</td><td align="center">[1, 32]</td><td align="center">float</td><td align="center">ND</td></tr>
41- <tr><td align="center">condition</td><td align="center">32</td><td align="center">bool</td><td align="center">ND</td></tr>41+ <tr><td align="center">condition</td><td align="center">[1, 32]</td><td align="center">bool</td><td align="center">ND</td></tr>
42- <tr><td rowspan="2" align="center">算子输出</td></tr>42+ <tr><td rowspan="2" align="center">样例输出</td></tr>
43- <tr><td align="center">dst</td><td align="center">32</td><td align="center">float</td><td align="center">ND</td></tr>43+ <tr><td align="center">dst</td><td align="center">[1, 32]</td><td align="center">float</td><td align="center">ND</td></tr>
44 44 
45 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">where_custom</td></tr>45 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">where_custom</td></tr>
46 </table>46 </table>
47 47 
48-- 算子实现： 48+- 样例实现：
49- 本样例中实现的是固定shape为输入src0[32]，src1[32]，condition[32]，输出dst[32]的where_custom算子。49+ 本样例中实现的是固定shape为输入src0[1, 32]，src1[1, 32]，condition[1, 32]，输出dst[1, 32]的where_custom样例。
50 50 
51- - Kernel实现 51+ - Kernel实现
52- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Where高阶API接口完成Where计算，得到最终结果，再搬出到外部存储上。
53 52 
54- where_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor src0Gm、src1Gm、conditionGm存储在src0Local、src1Local、conditionLocal中，Compute任务负责对src0Local、src1Local、conditionLocal执行Where计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。53+ 使用Where高阶API根据条件从两个源操作数中选择元素，支持张量和标量混合模式
55 54 
56 - 调用实现 55 - 调用实现
57 使用内核调用符<<<>>>调用核函数。56 使用内核调用符<<<>>>调用核函数。
58 57 
59## 编译运行 58## 编译运行
60 59 
61-在本样例根目录下执行如下步骤，编译并执行算子。60+在本样例根目录下执行如下步骤，编译并执行样例。
61+ 
62- 配置环境变量 62- 配置环境变量
63 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。63 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
64 - 默认路径，root用户安装CANN软件包64 - 默认路径，root用户安装CANN软件包
65+ 
65 ```bash66 ```bash
66 source /usr/local/Ascend/cann/set_env.sh67 source /usr/local/Ascend/cann/set_env.sh
67 ```68 ```
68 69 
69 - 默认路径，非root用户安装CANN软件包70 - 默认路径，非root用户安装CANN软件包
71+ 
70 ```bash72 ```bash
71 source $HOME/Ascend/cann/set_env.sh73 source $HOME/Ascend/cann/set_env.sh
72 ```74 ```
73 75 
74 - 指定路径install_path，安装CANN软件包76 - 指定路径install_path，安装CANN软件包
77+ 
75 ```bash78 ```bash
76 source ${install_path}/cann/set_env.sh79 source ${install_path}/cann/set_env.sh
77 ```80 ```
78- 81+ 
79- 样例执行82- 样例执行
83+ 
80 ```bash84 ```bash
81- mkdir -p build && cd build; # 创建并进入build目录85+ mkdir -p build && cd build; # 创建并进入build目录
82- cmake ..;make -j; # 编译工程86+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # 编译工程，默认npu模式
83 python3 ../scripts/gen_data.py # 生成测试输入数据87 python3 ../scripts/gen_data.py # 生成测试输入数据
84- ./demo # 执行编译生成的可执行程序，执行样例88+ ./demo # 执行编译生成的可执行程序，执行样例
85 ```89 ```
90+ 
91+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
92+
93+ 示例如下：
94+ 
95+ ```bash
96+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # cpu调试模式
97+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-3510 ..;make -j; # NPU仿真模式
98+ ```
99+ 
100+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
101+ 
102+- 编译选项说明
103+ 
104+ | 选项 | 可选值 | 说明 |
105+ |------|--------|------|
106+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
107+ | `CMAKE_ASC_ARCHITECTURES` | `dav-3510`（默认） | NPU 架构：dav-3510 对应 Ascend 950PR/Ascend 950DT |
108+ 
109+- 执行结果
110+ 
86 执行结果如下，说明精度对比成功。111 执行结果如下，说明精度对比成功。
112+ 
87 ```bash113 ```bash
88 test pass!114 test pass!
89- ```115+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/where/where.asc+24-2
@@ -11,13 +11,21 @@
11 11 
12/* !12/* !
13 * \file where.asc13 * \file where.asc
14- * \brief14+ * \brief 本样例基于Where高阶API实现条件选择功能，根据指定条件从两个源操作数中选择元素
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20 20 
21+#ifdef ASCENDC_CPU_DEBUG
22+#include "cpu_debug_launch.h"
23+#endif
24+ 
25+/**
26+ * @brief Where核函数实现类，演示Where API的使用场景
27+ * @tparam T 数据类型
28+ */
21template <typename T>29template <typename T>
22class KernelWhere {30class KernelWhere {
23public:31public:
@@ -67,6 +75,20 @@ public:
67 AscendC::LocalTensor<bool> conditionLocal = inQueueZ.DeQue<bool>();75 AscendC::LocalTensor<bool> conditionLocal = inQueueZ.DeQue<bool>();
68 AscendC::Duplicate(dstLocal, (T)0, shape);76 AscendC::Duplicate(dstLocal, (T)0, shape);
69 77 
78+ // 使用Where接口根据条件选择元素
79+ // 模板参数：
80+ // - T: 输入输出数据类型
81+ // 参数说明：
82+ // - dstLocal: 输出Tensor，存储选择结果
83+ // - src0Local: 第一个源操作数Tensor
84+ // - src1Local: 第二个源操作数Tensor
85+ // - conditionLocal: 条件Tensor，布尔类型
86+ // - dataSize: 计算元素个数
87+ // mode参数说明：
88+ // - 0: 张量对张量模式
89+ // - 1: src0为标量模式
90+ // - 2: src1为标量模式
91+ // - 3: src0和src1都为标量模式
70 if (mode == 0) {92 if (mode == 0) {
71 AscendC::Where<T>(dstLocal, src0Local, src1Local, conditionLocal, dataSize);93 AscendC::Where<T>(dstLocal, src0Local, src1Local, conditionLocal, dataSize);
72 } else if (mode == 1) {94 } else if (mode == 1) {
@@ -250,4 +272,4 @@ int32_t main(int32_t argc, char* argv[])
250 aclFinalize();272 aclFinalize();
251 273 
252 return 0;274 return 0;
253-}275+}
Mexamples/01_simd_cpp_api/03_libraries/12_math/xor/CMakeLists.txt+7-10
@@ -11,6 +11,9 @@
11 11 
12cmake_minimum_required(VERSION 3.16)12cmake_minimum_required(VERSION 3.16)
13 13 
14+set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
15+set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
16+ 
14find_package(ASC REQUIRED)17find_package(ASC REQUIRED)
15 18 
16project(kernel_samples LANGUAGES ASC CXX)19project(kernel_samples LANGUAGES ASC CXX)
@@ -19,21 +22,15 @@ add_executable(demo
19 xor.asc22 xor.asc
20)23)
21 24 
22-target_link_libraries(demo PRIVATE25+target_link_libraries(demo PRIVATE
23 tiling_api26 tiling_api
24 register27 register
25 platform28 platform
26 m29 m
27 dl30 dl
31+ graph_base
28)32)
29 33 
30-# ======================================================================================
31-# NPU 编译选项配置
32-#
33-# 说明：
34-# - 需根据实际部署的 NPU 硬件架构选择对应的 `npu-arch` 参数。
35-# ======================================================================================
36target_compile_options(demo PRIVATE34target_compile_options(demo PRIVATE
37- $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-2201>35+ $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=${CMAKE_ASC_ARCHITECTURES}>
38- # $<$<COMPILE_LANGUAGE:ASC>:--npu-arch=dav-3510>36+)
39-)
Mexamples/01_simd_cpp_api/03_libraries/12_math/xor/README.md+54-24
@@ -2,7 +2,7 @@
2 2 
3## 概述3## 概述
4 4 
5-本样例演示了基于Xor高阶API的算子实现。样例按元素执行Xor运算。5+本样例基于Xor高阶API实现按元素异或的功能。
6 6 
7## 支持的产品7## 支持的产品
8 8 
@@ -12,18 +12,18 @@
12 12 
13## 目录结构介绍13## 目录结构介绍
14 14 
15-```15+```plain
16├── xor16├── xor
17│ ├── scripts17│ ├── scripts
18-│ │ ├── gen_data.py // 输入数据和真值数据生成脚本18+│ │ └── gen_data.py // 输入数据和真值数据生成脚本
19│ ├── CMakeLists.txt // 编译工程文件19│ ├── CMakeLists.txt // 编译工程文件
20│ ├── data_utils.h // 数据读入写出函数20│ ├── data_utils.h // 数据读入写出函数
21-│ └── xor.asc // Ascend C算子实现 & 调用样例21+│ └── xor.asc // Ascend C样例实现 & 调用样例
22```22```
23 23 
24-## 算子描述24+## 样例描述
25 25 
26-- 算子功能： 26+- 样例功能：
27 按元素执行Xor运算，Xor（异或）的概念和运算规则如下： 27 按元素执行Xor运算，Xor（异或）的概念和运算规则如下：
28 概念：参加运算的两个数据，按二进制位进行“异或”运算。 28 概念：参加运算的两个数据，按二进制位进行“异或”运算。
29 运算规则：0^0=0；0^1=1；1^0=1；1^1=0；即：参加运算的两个对象，如果两个相应位为“异”（值不同），则该位结果为1，否则为 0【同0异1】。29 运算规则：0^0=0；0^1=1；1^0=1；1^1=0；即：参加运算的两个对象，如果两个相应位为“异”（值不同），则该位结果为1，否则为 0【同0异1】。
@@ -36,59 +36,89 @@
36 Xor(x, y) = (x \mid y) \& (\sim(x \& y))36 Xor(x, y) = (x \mid y) \& (\sim(x \& y))
37 $$37 $$
38 38 
39-- 算子规格： 39+- 样例规格：
40 <table>40 <table>
41- <tr><td rowspan="1" align="center">算子类型(OpType)</td><td colspan="4" align="center"> xor </td></tr>41+ <tr><td rowspan="1" align="center">样例类型(OpType)</td><td colspan="4" align="center"> xor </td></tr>
42 42 
43- <tr><td rowspan="4" align="center">算子输入</td></tr>43+ <tr><td rowspan="4" align="center">样例输入</td></tr>
44 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>44 <tr><td align="center">name</td><td align="center">shape</td><td align="center">data type</td><td align="center">format</td></tr>
45- <tr><td align="center">src0</td><td align="center">1024</td><td align="center">int16_t</td><td align="center">ND</td></tr>45+ <tr><td align="center">src0</td><td align="center">[1, 1024]</td><td align="center">int16_t</td><td align="center">ND</td></tr>
46- <tr><td align="center">src1</td><td align="center">1024</td><td align="center">int16_t</td><td align="center">ND</td></tr>46+ <tr><td align="center">src1</td><td align="center">[1, 1024]</td><td align="center">int16_t</td><td align="center">ND</td></tr>
47- <tr><td rowspan="2" align="center">算子输出</td></tr>47+ <tr><td rowspan="2" align="center">样例输出</td></tr>
48- <tr><td align="center">dst</td><td align="center">1024</td><td align="center">int16_t</td><td align="center">ND</td></tr>48+ <tr><td align="center">dst</td><td align="center">[1, 1024]</td><td align="center">int16_t</td><td align="center">ND</td></tr>
49 49 
50 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">xor_custom</td></tr>50 <tr><td rowspan="1" align="center">核函数名</td><td colspan="4" align="center">xor_custom</td></tr>
51 </table>51 </table>
52 52 
53-- 算子实现： 53+- 样例实现：
54- 本样例中实现的是固定shape为输入src0[1024]、是src1[1024]，输出dst[1024]的xor_custom算子。54+ 本样例中实现的是固定shape为输入src0[1, 1024]、src1[1, 1024]，输出dst[1, 1024]的xor_custom样例。
55 55 
56- - Kernel实现 56+ - Kernel实现
57- 计算逻辑是：Ascend C提供的矢量计算接口的操作元素都为LocalTensor，输入数据需要先搬运进片上存储，然后使用Xor高阶API接口完成Xor计算，得到最终结果，再搬出到外部存储上。
58 57 
59- xor_custom算子的实现流程分为3个基本任务：CopyIn，Compute，CopyOut。CopyIn任务负责将Global Memory上的输入Tensor src0Gm、src1Gm存储在src0Local、src1Local中，Compute任务负责对src0Local、src1Local执行Xor计算，计算结果存储在dstLocal中，CopyOut任务负责将输出数据从dstLocal搬运至Global Memory上的输出Tensor dstGm。58+ 使用Xor高阶API按元素进行异或运算，可选择使用临时buffer和指定计算元素个数
59+ 
60+ - Tiling实现
61+ 
62+ Host侧通过GetXorMaxMinTmpSize获取Xor接口计算所需的最大和最小临时空间。
60 63 
61 - 调用实现 64 - 调用实现
62 使用内核调用符<<<>>>调用核函数。65 使用内核调用符<<<>>>调用核函数。
63 66 
64## 编译运行 67## 编译运行
65 68 
66-在本样例根目录下执行如下步骤，编译并执行算子。69+在本样例根目录下执行如下步骤，编译并执行样例。
70+ 
67- 配置环境变量 71- 配置环境变量
68 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。72 请根据当前环境上CANN开发套件包的[安装方式](../../../../../docs/quick_start.md#prepare&install)，选择对应配置环境变量的命令。
69 - 默认路径，root用户安装CANN软件包73 - 默认路径，root用户安装CANN软件包
74+ 
70 ```bash75 ```bash
71 source /usr/local/Ascend/cann/set_env.sh76 source /usr/local/Ascend/cann/set_env.sh
72 ```77 ```
73 78 
74 - 默认路径，非root用户安装CANN软件包79 - 默认路径，非root用户安装CANN软件包
80+ 
75 ```bash81 ```bash
76 source $HOME/Ascend/cann/set_env.sh82 source $HOME/Ascend/cann/set_env.sh
77 ```83 ```
78 84 
79 - 指定路径install_path，安装CANN软件包85 - 指定路径install_path，安装CANN软件包
86+ 
80 ```bash87 ```bash
81 source ${install_path}/cann/set_env.sh88 source ${install_path}/cann/set_env.sh
82 ```89 ```
83- 90+ 
84- 样例执行91- 样例执行
92+ 
85 ```bash93 ```bash
86- mkdir -p build && cd build; # 创建并进入build目录94+ mkdir -p build && cd build; # 创建并进入build目录
87- cmake ..;make -j; # 编译工程95+ cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # 编译工程，默认npu模式
88 python3 ../scripts/gen_data.py # 生成测试输入数据96 python3 ../scripts/gen_data.py # 生成测试输入数据
89- ./demo # 执行编译生成的可执行程序，执行样例97+ ./demo # 执行编译生成的可执行程序，执行样例
90 ```98 ```
99+ 
100+ 使用 CPU调试 或 NPU仿真 模式时，添加 `-DCMAKE_ASC_RUN_MODE=cpu` 或 `-DCMAKE_ASC_RUN_MODE=sim` 参数即可。
101+
102+ 示例如下：
103+ 
104+ ```bash
105+ cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # cpu调试模式
106+ cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..;make -j; # NPU仿真模式
107+ ```
108+ 
109+ > **注意：** 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 `rm CMakeCache.txt` 后重新 cmake。
110+ 
111+- 编译选项说明
112+ 
113+ | 选项 | 可选值 | 说明 |
114+ |------|--------|------|
115+ | `CMAKE_ASC_RUN_MODE` | `npu`（默认）、`cpu`、`sim` | 运行模式：NPU 运行、CPU调试、NPU仿真 |
116+ | `CMAKE_ASC_ARCHITECTURES` | `dav-2201`（默认）、`dav-3510` | NPU 架构：dav-2201 对应 Atlas A2/A3 系列，dav-3510 对应 Ascend 950PR/Ascend 950DT |
117+ 
118+- 执行结果
119+ 
91 执行结果如下，说明精度对比成功。120 执行结果如下，说明精度对比成功。
121+ 
92 ```bash122 ```bash
93 test pass!123 test pass!
94- ```124+ ```
Mexamples/01_simd_cpp_api/03_libraries/12_math/xor/xor.asc+28-5
@@ -11,13 +11,22 @@
11 11 
12/* !12/* !
13 * \file xor.asc13 * \file xor.asc
14- * \brief14+ * \brief 本样例基于Xor高阶API实现按位异或运算功能，按元素执行异或运算
15 */15 */
16 16 
17#include "acl/acl.h"17#include "acl/acl.h"
18#include "data_utils.h"18#include "data_utils.h"
19#include "kernel_operator.h"19#include "kernel_operator.h"
20+#include "tiling/tiling_api.h"
20 21 
22+#ifdef ASCENDC_CPU_DEBUG
23+#include "cpu_debug_launch.h"
24+#endif
25+ 
26+/**
27+ * @brief Xor核函数实现类，演示Xor API的使用场景
28+ * @tparam T 数据类型
29+ */
21template <typename T>30template <typename T>
22class KernelXor {31class KernelXor {
23public:32public:
@@ -64,6 +73,16 @@ public:
64 getTempBuffer = tempBuffer.Get<uint8_t>();73 getTempBuffer = tempBuffer.Get<uint8_t>();
65 }74 }
66 75 
76+ // 使用Xor接口按元素进行异或运算
77+ // 模板参数：
78+ // - T: 输入输出数据类型
79+ // - false: 是否复用源操作数
80+ // 参数说明：
81+ // - dstLocal: 输出Tensor，存储计算结果
82+ // - src0Local: 第一个输入Tensor
83+ // - src1Local: 第二个输入Tensor
84+ // - getTempBuffer: 临时buffer，用于提高精度
85+ // - calCount: 计算元素个数
67 if ((tmpBufSize > 0) && (calCount > 0)) {86 if ((tmpBufSize > 0) && (calCount > 0)) {
68 AscendC::Xor<T, false>(dstLocal, src0Local, src1Local, getTempBuffer, calCount);87 AscendC::Xor<T, false>(dstLocal, src0Local, src1Local, getTempBuffer, calCount);
69 } else if (tmpBufSize > 0) {88 } else if (tmpBufSize > 0) {
@@ -97,11 +116,10 @@ private:
97 uint32_t dataSize = 0;116 uint32_t dataSize = 0;
98};117};
99 118 
100-__global__ __vector__ void xor_custom(GM_ADDR srcGm, GM_ADDR src1Gm, GM_ADDR dstGm)119+__global__ __vector__ void xor_custom(GM_ADDR srcGm, GM_ADDR src1Gm, GM_ADDR dstGm, uint32_t tmpBufSize)
101{120{
102 AscendC::TPipe pipe;121 AscendC::TPipe pipe;
103 constexpr uint32_t srcSize = 1024;122 constexpr uint32_t srcSize = 1024;
104- constexpr uint32_t tmpBufSize = 128;
105 constexpr uint32_t calCount = 1024;123 constexpr uint32_t calCount = 1024;
106 constexpr uint32_t apiMode = 0;124 constexpr uint32_t apiMode = 0;
107 KernelXor<int16_t> op;125 KernelXor<int16_t> op;
@@ -150,6 +168,11 @@ int32_t main(int32_t argc, char* argv[])
150 size_t param3FileSize = 1024 * sizeof(int16_t);168 size_t param3FileSize = 1024 * sizeof(int16_t);
151 uint32_t numBlocks = 1;169 uint32_t numBlocks = 1;
152 170 
171+ ge::Shape shape{{1024}};
172+ uint32_t maxValue = 0;
173+ uint32_t minValue = 0;
174+ AscendC::GetXorMaxMinTmpSize(shape, sizeof(int16_t), false, maxValue, minValue);
175+ 
153 aclInit(nullptr);176 aclInit(nullptr);
154 aclrtContext context;177 aclrtContext context;
155 int32_t deviceId = 0;178 int32_t deviceId = 0;
@@ -177,7 +200,7 @@ int32_t main(int32_t argc, char* argv[])
177 aclrtMallocHost((void**)(&param3Host), param3FileSize);200 aclrtMallocHost((void**)(&param3Host), param3FileSize);
178 aclrtMalloc((void**)&param3Device, param3FileSize, ACL_MEM_MALLOC_HUGE_FIRST);201 aclrtMalloc((void**)&param3Device, param3FileSize, ACL_MEM_MALLOC_HUGE_FIRST);
179 202 
180- xor_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device);203+ xor_custom<<<numBlocks, nullptr, stream>>>(param1Device, param2Device, param3Device, minValue);
181 aclrtSynchronizeStream(stream);204 aclrtSynchronizeStream(stream);
182 205 
183 aclrtFree(param1Device);206 aclrtFree(param1Device);
@@ -205,4 +228,4 @@ int32_t main(int32_t argc, char* argv[])
205 aclFinalize();228 aclFinalize();
206 229 
207 return 0;230 return 0;
208-}231+}