Ascend C Swish/Silu 激活算子样例全解析：从高阶 API 调用到编译运行验证（asc-devkit 实战）

导读

本篇文章以 CANN asc-devkit 开源仓库中的 Swish 样例（examples/01_simd_cpp_api/04_advanced_api/01_activation/swish/README.md）为核心，完整讲解基于 Ascend C 高阶 API（Swish / Silu）在昇腾 AI 处理器上实现逐元素激活算子的开发流程。读完本文，你将掌握：Swish 与 Silu 的数学关系与换算方式、高阶 API 的 Kernel 端调用范式、GetSwishTmpSize / GetSiluTmpSize 的 Tiling 用法、编译宏 USE_SILU_MODE 控制的双模式构建流程，以及 CPU 调试 / NPU 仿真 / NPU 实际运行三种运行模式与精度对比机制。文中所有结论均可在仓库源码中找到对应实现。

一、样例概述与数学原理

1.1 Swish 与 Silu 的数学关系

本样例面向深度学习模型的激活函数场景，基于 Swish/Silu 高阶 API 实现激活函数计算。两个 API 具有紧密的数学关系：

• Swish（默认模式）：y = x / (1 + exp(-beta * x))，其中 beta 为可调节参数。本样例默认取 beta = 1.702（GELU 近似值）。
• Silu（Silu 模式）：Swish 在 beta = 1 时的特例，y = x / (1 + exp(-x))。相比 ReLU，Silu 曲线更平滑，负半轴不硬截断，梯度不会完全消失。

从实现角度，Swish 也可写作 Swish(x, beta) = x * sigmoid(beta * x)，Silu 即 Silu(x) = x * sigmoid(x)。当 beta = 1 时二者完全等价，可互相替换。样例源码 swish.asc 的注释也明确说明了这一关系，并通过编译宏 USE_SILU_MODE 在两种模式间切换，输入数据类型为 float，输入 Tensor 元素个数为 32。

样例的计算公式如下：

$$dstLocal_i = Swish(srcLocal_i) = \frac{srcLocal_i}{1 + e^{-\beta \cdot srcLocal_i}}$$

$$dstLocal_i = Silu(srcLocal_i) = \frac{srcLocal_i}{1 + e^{-srcLocal_i}}$$

1.2 样例规格

本样例实现的是固定 shape 的 swish_custom 核函数：输入 src[1, 32]、输出 dst[1, 32]，数据类型 float、Format 为 ND，单核逐元素计算。具体规格如下表：

项目	规格
样例类型（OpType）	swish / silu
样例输入 src	shape [1, 32]，data type float，format ND
样例输出 dst	shape [1, 32]，data type float，format ND
核函数名	swish_custom

从 swish.asc 的 Kernel 入口函数可以看到，inputSize 被定义为 constexpr uint32_t inputSize = 32，与规格表完全对应：

__global__ __vector__ void swish_custom(__gm__ uint8_t* srcGm, __gm__ uint8_t* dstGm)
{
    AscendC::TPipe pipe;
    constexpr uint32_t inputSize = 32;
#ifdef USE_SILU_MODE
    // Silu模式：beta=1，Swish的特例
    float betaValue = 1.0f;
#else
    // Swish模式：beta=1.702，GELU近似
    float betaValue = 1.702f;
#endif
    KernelSwish<float> op;
    op.Init(srcGm, dstGm, inputSize, betaValue, &pipe);
    op.Process();
}

betaValue 由编译宏决定：Swish 模式取 1.702f，Silu 模式取 1.0f，随后作为参数传入算子类完成初始化。

1.3 支持的产品及 CANN 软件版本

根据样例 README，本样例支持的产品与 CANN 版本对应关系如下：

产品	CANN 软件版本
Ascend 950PR&950DT 系列产品	>= CANN 9.1.0
Atlas A3 系列产品	>= CANN 9.0.0
Atlas A2 系列产品	>= CANN 9.0.0

在构建参数上，Atlas A2/A3 系列对应 dav-2201 架构，Ascend 950PR&950DT 系列对应 dav-3510 架构（详见下文编译选项说明）。

二、目录结构与文件角色

样例根目录 examples/01_simd_cpp_api/04_advanced_api/01_activation/swish 下包含以下文件：

├── swish
│   ├── scripts
│   │   ├── gen_data.py         // 输入数据和真值数据生成脚本（支持Swish/Silu两种模式）
│   ├── CMakeLists.txt          // 编译工程文件（支持 -DUSE_SILU_MODE）
│   ├── data_utils.h            // 数据读入写出函数
│   ├── swish.asc               // Ascend C算子实现 & 调用样例（两种模式合并）
│   └── README.md               // 样例说明文档

各文件职责如下：

• swish.asc：核心源码，包含 Device 端 Kernel 实现（KernelSwish 算子类与 swish_custom 入口函数）和 Host 端调用逻辑（ACL 环境初始化、内存分配、核函数调用、精度对比）。
• CMakeLists.txt：编译工程文件，通过 option(USE_SILU_MODE ...) 声明编译宏，并在 USE_SILU_MODE=ON 时向 demo 目标注入 USE_SILU_MODE 宏定义。
• scripts/gen_data.py：生成二进制输入文件 input/input_src.bin 与真值文件 output/golden.bin，支持 --silu-mode 参数。
• data_utils.h：提供 ReadFile / WriteFile 两个二进制文件读写工具函数，供 Host 端加载输入、落盘输出。

三、Kernel 端实现剖析

3.1 算子类的三段式流程：CopyIn → Compute → CopyOut

KernelSwish 算子类（模板类型 T，本样例实例化为 float）在 Init 阶段完成三件事：

1. 通过 SetGlobalBuffer 将全局内存 srcGm、dstGm 绑定到 GlobalTensor<T>；
2. 保存 betaValue 与 dataSize；
3. 通过 pipe->InitBuffer 为输入队列 inQueueX 和输出队列 outQueue 各申请一块大小为 dataSize * sizeof(T) 的本地内存（UB）缓冲区。

Process() 则按经典流水线顺序执行：

__aicore__ inline void Process()
{
    CopyIn();
    Compute();
    CopyOut();
}

各阶段实现如下：

• CopyIn：从输入队列分配 LocalTensor，调用 AscendC::DataCopy(srcLocal, srcGlobal, dataSize) 将全局内存数据搬入 UB，再 EnQue 入队；
• Compute：DeQue 取出输入、AllocTensor 分配输出，按编译宏分支调用高阶 API 完成计算，随后 EnQue 输出并 FreeTensor 释放输入；
• CopyOut：DeQue 输出 Tensor，调用 AscendC::DataCopy(dstGlobal, dstLocal, dataSize) 将结果搬回全局内存。

3.2 高阶 API 调用方式

计算阶段是样例的核心。两种模式的调用如下：

Swish 模式（默认）：

AscendC::Swish(dstLocal, srcLocal, dataSize, scalarValue);

Silu 模式：

AscendC::Silu<T, false>(dstLocal, srcLocal, dataSize);

对应源码中的分支逻辑（swish.asc）：

#ifdef USE_SILU_MODE
    // Silu模式：等价于 Swish(beta=1)
    // 以下两种调用结果完全等价：
    //   AscendC::Silu<T, false>(dstLocal, srcLocal, dataSize);
    //   AscendC::Swish(dstLocal, srcLocal, dataSize, 1.0f);
    AscendC::Silu<T, false>(dstLocal, srcLocal, dataSize);
#else
    // Swish模式：beta=1.702（GELU近似值）
    AscendC::Swish(dstLocal, srcLocal, dataSize, betaValue);
#endif

公开接口的签名可以在 include/adv_api/activation/swish.h 与 include/adv_api/activation/silu.h 中看到：

• Swish：template <typename T, bool isReuseSource = false> void Swish(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, uint32_t dataSize, const T scalarValue)，其中 scalarValue 即公式中的 β 参数；
• Silu：template <typename T, bool isReuseSource = false> void Silu(const LocalTensor<T>& dstLocal, const LocalTensor<T>& srcLocal, uint32_t dataSize)。

两个接口都声明支持 half 和 float 数据类型，且仅在 AI Vector Core 上生效（if ASCEND_IS_AIC 直接返回），并明确禁止在 --enable-simt（纯 SIMT）编译模式下使用。isReuseSource 模板参数用于指示是否复用源操作数输入空间。

3.3 底层计算原理：四步指令流水线

高阶 API 并非黑盒，其底层通过基础 API 组合实现。查看通用实现 impl/adv_api/detail/activation/swish/swish_common_impl.h，SwishCalcSimplified 将 swish(x) = x / (1 + e^(-βx)) 拆解为四条向量指令：

// swish(x) = x / (1 + e^(-βx))
// x1 = 1 + e^(-βx)
Muls<T, false>(dstAddr, srcAddr, scalarValue, MASK_PLACEHOLDER, repeatTimes, unaryParams); // dst = -βx
PipeBarrier<PIPE_V>();
Exp<T, false>(dstAddr, dstAddr, MASK_PLACEHOLDER, repeatTimes, unaryParams);              // dst = e^(-βx)
PipeBarrier<PIPE_V>();
Adds<T, false>(dstAddr, dstAddr, static_cast<T>(1), MASK_PLACEHOLDER, repeatTimes, unaryParams); // dst = 1 + e^(-βx)
PipeBarrier<PIPE_V>();
Div<T, false>(dstAddr, srcAddr, dstAddr, MASK_PLACEHOLDER, repeatTimes, binaryParams);     // dst = x / (1 + e^(-βx))
PipeBarrier<PIPE_V>();

注意：SwishCompute 在调用前会将 scalarValue 取负（scalar = -1 * scalarValue），因此第一条 Muls 实际计算的是 -βx，随后 Exp → Adds(+1) → Div(src, dst) 依次完成指数、加 1、相除。每步之间通过 PipeBarrier<PIPE_V>() 插入向量流水线屏障，保证指令顺序正确。

Silu 的实现（impl/adv_api/detail/activation/silu/silu_common_impl.h）走同一条流水线，只是把 β 固定为 1：Muls(-1) → Exp → Adds(1) → Div(src, dst)。

在非 2201/2002/5102/5162 架构上，SwishCompute / SiluCompute 还会针对数据量做分块处理：根据 MAX_REPEAT_*_SIZE、ONE_REPEAT_*_SIZE 与 MAX_REPEAT_TIMES 计算 repeat 轮数，并用 SetVectorMask 处理尾部（tail）数据，保证任意元素个数都能正确计算。这正是样例将数据量设为 32 也能安全运行的原因。

3.4 不同架构的实现差异

高阶 API 在不同 NPU 架构下会引入不同的实现头文件。swish.h 中通过 __NPU_ARCH__ 宏判断：当架构为 3510 / 9201 / 5102 / 5162 / 3003 / 3113 时包含 swish_3510_impl.h，否则包含 swish_common_impl.h；silu.h 同样如此。

以 swish_3510_impl.h 为例，它在 3510 架构上改用寄存器级 SIMD 实现：通过 Reg::LoadAlign 加载数据、Reg::Muls / Reg::Exp / Reg::Adds / Reg::Div 完成四步计算、Reg::StoreAlign 存储结果，并用 Reg::UpdateMask<T>(count) 逐 repeat 更新掩码。该实现还通过 static_assert(SupportType<T, half, float>(), ...) 在编译期强制数据类型为 half/float，并通过 CheckTensorPosition、CheckCalCount 检查 Tensor 位置（要求 VECIN/VECOUT/VECCALC）与计算元素个数。可以看到，同一套高阶 API 在不同架构上由编译器自动选择优化实现，开发者无需关心底层差异。

四、Host 端与 Tiling 实现

4.1 Tiling：GetSwishTmpSize / GetSiluTmpSize

本样例为单核逐元素计算场景，无复杂分核逻辑。Host 端在调用核函数前，通过 Tiling 接口获取高阶 API 所需的临时缓冲区大小：

std::vector<int64_t> shapeVec = {inputSize};
AscendC::TensorShape srcShape(shapeVec);
uint32_t maxValue = 0;
uint32_t minValue = 0;
#ifdef USE_SILU_MODE
    AscendC::GetSiluTmpSize(srcShape, sizeof(float), false, maxValue, minValue);
#else
    AscendC::GetSwishTmpSize(srcShape, sizeof(float), false, maxValue, minValue);
#endif

接口签名定义在 include/adv_api/activation/swish_tiling.h 与 include/adv_api/activation/silu_tiling.h：传入 srcShape（输入 Tensor shape）、typeSize（dtype 字节数）、isReuseSource（是否复用源操作数空间），通过引用参数 maxValue / minValue 返回临时本地空间的最大/最小申请量，开发者可据此区间选择合适的空间大小作为 Tiling 参数。

从当前仓库实现看（impl/adv_api/tiling/activation/swish_tiling.cpp 与 silu_tiling.cpp），GetSwishTmpSize / GetSiluTmpSize 将 maxValue 与 minValue 均置为 0——结合 3.3 节可知，Swish/Silu 的中间计算全部在 dstLocal 上就地完成，无需额外的临时 UB 空间，因此样例中并未据此分配额外缓冲区。这也解释了为什么样例的 Kernel 端只需输入、输出两块缓冲即可运行。

4.2 ACL 环境初始化与内存管理

main 函数首先完成 ACL（Ascend Computing Language）运行环境初始化：

aclInit(nullptr);
aclrtContext context;
aclrtSetDevice(deviceId);              // deviceId = 0
aclrtCreateContext(&context, deviceId);
aclrtStream stream = nullptr;
aclrtCreateStream(&stream);

随后按 inputSize * sizeof(float)（即 32 × 4 = 128 字节）分别申请 Host 与 Device 内存：

• 输入：aclrtMallocHost 申请 Host 侧 srcHost，aclrtMalloc(..., ACL_MEM_MALLOC_HUGE_FIRST) 申请 Device 侧 srcDevice；
• 输出：同样申请 dstHost / dstDevice；
• 数据搬运：ReadFile 读取 ./input/input_src.bin 到 srcHost，再 aclrtMemcpy(..., ACL_MEMCPY_HOST_TO_DEVICE) 拷入 Device。

4.3 核函数调用与精度对比

Host 端使用内核调用符 <<<>>> 启动核函数：

swish_custom<<<numBlocks, 0, stream>>>(srcDevice, dstDevice);
aclrtSynchronizeStream(stream);

numBlocks 为 1，即单核执行。执行完成后将结果从 Device 拷回 Host，写出 ./output/output.bin，并调用 CompareResult 与真值文件 ./output/golden.bin 逐元素对比。对比逻辑（swish.asc）采用绝对误差与相对误差双重判据，阈值 EPS = 1e-4：

constexpr float EPS = 1e-4;
// 当绝对误差与相对误差同时超过 EPS 时判为错误
float ae = std::abs(a - b);
float re = ae / std::abs(b);
if (ae > EPS && re > EPS) { ... }

对比通过后打印 CompareResult golden.bin success!，最终 main 返回 0 并输出 test pass!。程序退出前依次释放 Device/Host 内存、销毁 Stream 与 Context、aclrtResetDevice 并 aclFinalize()，资源管理完整规范。

真值数据由 scripts/gen_data.py 生成：输入 src 取 np.random.uniform(-4.0, 4.0, [32]) 的 float32 随机数，真值分别按 swish(x, beta=1.702) 或 silu(x) 公式用 NumPy 计算后以二进制格式写入 ./output/golden.bin。输入数据限定在 [-4, 4] 区间，可避免指数运算在极端输入下产生数值溢出，保证对比稳定。

五、编译运行与验证

5.1 环境准备

在样例根目录执行编译运行前，需先根据当前环境上 CANN 开发套件包的安装方式配置环境变量，详细步骤参考 docs/zh/quick_start.md#prepare&install。典型做法是执行：

source ${install_path}/cann/set_env.sh

说明： ${install_path} 为 CANN 包安装目录，未指定安装目录时默认安装至 /usr/local/Ascend 下。

5.2 Swish 模式编译运行（默认）

在样例根目录 examples/01_simd_cpp_api/04_advanced_api/01_activation/swish 下执行：

mkdir -p build && cd build;
cmake -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..; make -j;
python3 ../scripts/gen_data.py
./demo

执行成功输出：

test pass!

5.3 Silu 模式编译运行

Silu 模式需同时开启 USE_SILU_MODE 编译选项，并用 --silu-mode 让数据脚本按 Silu 公式生成真值：

mkdir -p build && cd build;
cmake -DUSE_SILU_MODE=ON -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..; make -j;
python3 ../scripts/gen_data.py --silu-mode
./demo

USE_SILU_MODE 宏会通过 CMakeLists.txt 中的 target_compile_definitions(demo PRIVATE USE_SILU_MODE) 注入到编译命令，从而在 swish.asc 中切换 Silu API 与 beta=1.0f。

5.4 CPU 调试与 NPU 仿真模式

无 NPU 硬件时，可通过 CMAKE_ASC_RUN_MODE 切换到 CPU 调试或 NPU 仿真模式：

cmake -DCMAKE_ASC_RUN_MODE=cpu -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..; make -j; # cpu调试模式
cmake -DCMAKE_ASC_RUN_MODE=sim -DCMAKE_ASC_ARCHITECTURES=dav-2201 ..; make -j; # NPU仿真模式

注意： 切换编译模式前需清理 cmake 缓存，可在 build 目录下执行 rm CMakeCache.txt 后重新 cmake。

5.5 编译选项说明

选项	可选值	说明
CMAKE_ASC_RUN_MODE	npu（默认）、cpu、sim	运行模式：NPU 运行、CPU 调试、NPU 仿真
CMAKE_ASC_ARCHITECTURES	dav-2201（默认）、dav-3510	NPU 架构：dav-2201 对应 Atlas A2 系列产品和 Atlas A3 系列产品，dav-3510 对应 Ascend 950PR&950DT 系列产品
USE_SILU_MODE	OFF（默认）、ON	样例模式：OFF 为 Swish，ON 为 Silu

在 CMakeLists.txt 中可以看到这些选项的默认值来源：

set(CMAKE_ASC_RUN_MODE "npu" CACHE STRING "Run mode: npu, cpu, sim")
set(CMAKE_ASC_ARCHITECTURES "dav-2201" CACHE STRING "NPU architecture: dav-2201, dav-3510")
option(USE_SILU_MODE "Enable Silu mode (default: OFF)" OFF)

此外，工程通过 find_package(ASC REQUIRED) 引入 Ascend C 编译环境，project(kernel_samples LANGUAGES ASC CXX) 声明使用 ASC 语言扩展，并链接了 tiling_api、register、platform、unified_dlog、graph_base 等运行依赖库。

5.6 执行结果与数据生成机制

./demo 运行完成后，会在 build 目录下生成 input/input_src.bin（输入）、output/golden.bin（真值）与 output/output.bin（算子输出）。精度对比成功时输出 test pass!。若对比失败，CompareResult 会打印首个错误位置处输出值与真值：

CompareResult golden.bin failed output is %lf, golden is %lf

六、仓库中的配套测试与扩展阅读

仓库在 tests/api/adv_api/activation/swish/test_operator_swish.cpp 中提供了 Swish 算子的单元测试：测试覆盖 half 与 float 两种数据类型，数据量从 32 到 32768 共 10 档，通过 SwishTestSuite 参数化测试验证不同 shape 下的计算正确性。其 Kernel 结构（CopyIn → Compute(Swish) → CopyOut）与样例 swish.asc 完全同构，可作为批量数据规模的参考实现。

关于高阶 API 的更多背景，可参考仓库文档 docs/zh/quick_start.md（环境准备与编译安装）以及本样例同目录下的 README_en.md（英文版说明）。如需在昇腾 AI 处理器上实现其他激活函数（如 GELU、Sigmoid、Softmax 等），可参考 impl/adv_api/tiling/activation 目录下对应的 Tiling 实现，其调用范式与本样例一致。