企业官网建设流程全解析

CANN ops-nn Sigmoid 算子深度指南：aclnn 两段式接口、AI Core 内核实现与图模式调用

【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn

本文以 CANN ops-nn 仓库中的 Sigmoid 算子的官方文档（activation/sigmoid/README.md 与 aclnnSigmoid&aclnnInplaceSigmoid.md）为主线，结合该算子模块的源码、样例与测试，系统讲解 Sigmoid 在 NPU 上的功能定义、产品支持范围、aclnnSigmoid/aclnnInplaceSigmoid两段式调用方法、底层 AI Core 内核实现以及算子 IR 构图方式。读完本文，你将能够独立完成 Sigmoid 算子在 Ascend 设备上的算子级开发与验证，并理解从 aclnn 接口到 Tiling、再到 Kernel 的完整执行链路。

一、功能说明与数学定义

Sigmoid 是神经网络中最常用的激活函数之一，将任意实数输入映射到 (0, 1) 开区间内，常用于二分类输出层、LSTM/GRU 等循环神经网络的门控单元以及注意力机制的归一化环节。在本仓库中，该算子位于activation/sigmoid目录，其功能为：对输入 Tensor 逐元素（element-wise）完成 Sigmoid 运算。

计算公式为：

$$ out = {\frac{1} {1+{e}^{-input}}} $$

该函数具有如下典型性质：

• 输出值域为 (0, 1)，单调递增；
• 当输入趋近正无穷时输出趋近 1，趋近负无穷时输出趋近 0；
• 关于原点中心对称：sigmoid(-x) = 1 - sigmoid(x)；
• 导数可用自身表示：sigmoid'(x) = sigmoid(x) * (1 - sigmoid(x))，这一性质使其在反向传播中计算高效，也是 activation/sigmoid_grad 等梯度算子实现的基础。

从算子原型看，Sigmoid 的输入输出形状完全一致（逐元素运算），由 op_graph/sigmoid_proto.h 中的 IR 定义可见，其输入x与输出y均使用TensorType::UnaryDataType()，且该 IR 明确标注“Compatible with the TensorFlow operator Sigmoid”，即与 TensorFlow 框架的 Sigmoid 算子语义对齐。

二、产品支持情况

根据官方文档，Sigmoid 算子在各产品上的支持情况如下：

注：以上为算子级支持情况；aclnnSigmoid/aclnnInplaceSigmoid接口级支持情况与上表一致，详见接口文档 aclnnSigmoid&aclnnInplaceSigmoid.md。

从源码可印证该支持范围：在 op_host/sigmoid_def.cpp 中，算子通过this->AICore().AddConfig("ascend950", aicoreConfig)与AddConfig("ascend350", aicoreConfig)注册了 ascend950、ascend350 两个平台的 AI Core 配置（对应文档中的 Ascend 950 系列与 Atlas A2/A3 系列等）；而 Atlas 200I/500 A2 推理产品不在注册列表中，因此标记为不支持。

三、参数说明

Sigmoid 算子仅有输入x和输出y两个 Tensor 参数，具体规格如下：

需要说明的是，上表为算子（Kernel）级支持的数据类型。在 aclnn 接口层，通过 op_api/aclnn_sigmoid.cpp 中的DTYPE_SUPPORT_LIST可以看到，接口输入self还额外支持DOUBLE、INT8、INT16、INT32、INT64、UINT8、BOOL、COMPLEX64、COMPLEX128等更宽泛的类型——这些类型会在接口层通过Cast算子统一转换后再送入 AI Core Kernel 计算，这也是“算子层只支持三种浮点类型、而接口层支持更多类型”的原因所在。

约束说明：官方文档明确标注“无”额外约束。在实现层面，唯一的形状约束是输入self与输出out的 shape 必须一致（见下文参数校验逻辑）。

四、调用方式总览

Sigmoid 算子在 CANN 中支持两种调用方式：

其中图模式对应的 IR 定义位于 op_graph/sigmoid_proto.h，通过REG_OP(Sigmoid)注册算子原型；Host 侧的形状推导在 op_host/sigmoid_infershape.cpp 中实现，直接复用Ops::Base::InferShape4Elewise（逐元素通用推断逻辑），保证输出 shape 与输入一致。

五、两段式 aclnn 接口详解

acclnSigmoid与acclnInplaceSigmoid实现相同的计算功能，区别仅在于结果存储方式：

• aclnnSigmoid：需新建一个输出张量对象out存储计算结果，输入输出分离；
• aclnnInplaceSigmoid：无需新建输出张量对象，直接在输入张量的内存中覆写计算结果（in-place 操作），节省一倍显存。

与 CANN 其他算子一样，这两个接口均采用两段式调用：必须先调用第一段xxxGetWorkspaceSize接口获取 workspace 大小并完成入参校验、构建aclOpExecutor，再调用第二段xxx接口在指定 Stream 上真正下发执行。两段式接口的通用原理可参考 docs/zh/context/basic_concept.md 中的 aclnn 两段式 API 说明。

5.1 函数原型

aclnnStatus aclnnSigmoidGetWorkspaceSize( const aclTensor *self, aclTensor *out, uint64_t *workspaceSize, aclOpExecutor **executor) aclnnStatus aclnnSigmoid( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream) aclnnStatus aclnnInplaceSigmoidGetWorkspaceSize( aclTensor* selfRef, uint64_t* workspaceSize, aclOpExecutor** executor) aclnnStatus aclnnInplaceSigmoid( void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, const aclrtStream stream)

5.2 aclnnSigmoidGetWorkspaceSize 参数说明

5.3 aclnnInplaceSigmoidGetWorkspaceSize 参数说明

注意：in-place 接口在部分老产品（Atlas 推理系列产品、Atlas 训练系列产品）上仅支持FLOAT、FLOAT16、DOUBLE、COMPLEX64、COMPLEX128，不支持BFLOAT16。这一点在源码 op_api/aclnn_sigmoid.cpp 中有对应逻辑：GetSelfRefDtypeList()会根据当前 NPU 架构（DAV_2201或 RegBase 架构）动态返回不同的 in-place 类型支持列表，BF16 仅在 Ascend 950 及 Atlas A2/A3 系列等新架构上开放。

5.4 两段接口（aclnnSigmoid / aclnnInplaceSigmoid）公共参数

5.5 返回值与错误码

两段接口均返回aclnnStatus状态码。第一段接口会完成入参校验，出现以下场景时报错：

aclnnSigmoidGetWorkspaceSize 校验错误：

aclnnInplaceSigmoidGetWorkspaceSize 校验错误：

这些校验在源码中的落点是 op_api/aclnn_sigmoid.cpp 的CheckParams/CheckInplaceParams函数：依次检查空指针（CheckNotNull）、数据类型与 BF16 平台兼容性（CheckDtypeValid，其中OP_CHECK_DTYPE_NOT_SUPPORT映射到ACLNN_ERR_PARAM_INVALID161002）、shape 一致性（CheckShape，要求 self 与 out 的 shape 完全一致且维度不超过 8）。

5.6 约束说明

• 确定性计算：aclnnSigmoid与aclnnInplaceSigmoid默认即为确定性实现，同一输入多次执行结果一致，无需额外配置。

六、完整调用示例（可直接运行）

下面以aclnnSigmoid为例给出完整的调用流程，该示例可在 examples/test_aclnn_sigmoid.cpp 中找到完整可编译源码；aclnnInplaceSigmoid的完整示例见 examples/test_aclnn_inplace_sigmoid.cpp。编译与运行环境准备可参考仓库文档中的样例编译说明。

6.1 公共工具代码

#include <cinttypes> #include <iostream> #include <vector> #include "acl/acl.h" #include "aclnnop/aclnn_sigmoid.h" #define CHECK_RET(cond, return_expr) \ do { \ if (!(cond)) { \ return_expr; \ } \ } while (0) #define LOG_PRINT(message, ...) \ do { \ printf(message, ##__VA_ARGS__); \ } while (0) int64_t GetShapeSize(const std::vector<int64_t>& shape) { int64_t shapeSize = 1; for (auto i : shape) { shapeSize *= i; } return shapeSize; } int Init(int32_t deviceId, aclrtStream* stream) { // 固定写法，资源初始化 auto ret = aclInit(nullptr); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclInit failed. ERROR: %d\n", ret); return ret); ret = aclrtSetDevice(deviceId); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSetDevice failed. ERROR: %d\n", ret); return ret); ret = aclrtCreateStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtCreateStream failed. ERROR: %d\n", ret); return ret); return 0; } template <typename T> int CreateAclTensor(const std::vector<T>& hostData, const std::vector<int64_t>& shape, void** deviceAddr, aclDataType dataType, aclTensor** tensor) { auto size = GetShapeSize(shape) * sizeof(T); // 调用aclrtMalloc申请device侧内存 auto ret = aclrtMalloc(deviceAddr, size, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMalloc failed. ERROR: %d\n", ret); return ret); // 调用aclrtMemcpy将host侧数据拷贝到device侧内存上 ret = aclrtMemcpy(*deviceAddr, size, hostData.data(), size, ACL_MEMCPY_HOST_TO_DEVICE); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtMemcpy failed. ERROR: %d\n", ret); return ret); // 计算连续tensor的strides std::vector<int64_t> strides(shape.size(), 1); for (int64_t i = shape.size() - 2; i >= 0; i--) { strides[i] = shape[i + 1] * strides[i + 1]; } // 调用aclCreateTensor接口创建aclTensor（ND格式，storageOffset为0） *tensor = aclCreateTensor(shape.data(), shape.size(), dataType, strides.data(), 0, aclFormat::ACL_FORMAT_ND, shape.data(), shape.size(), *deviceAddr); return 0; }

6.2 aclnnSigmoid 主流程（非原地版）

int main() { // 1. （固定写法）device/stream初始化，参考acl API；deviceId根据实际设备填写 int32_t deviceId = 0; aclrtStream stream; auto ret = Init(deviceId, &stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("Init acl failed. ERROR: %d\n", ret); return ret); // 2. 构造输入与输出 std::vector<int64_t> selfShape = {2, 2}; std::vector<int64_t> outShape = {2, 2}; void* selfDeviceAddr = nullptr; void* outDeviceAddr = nullptr; aclTensor* self = nullptr; aclTensor* out = nullptr; std::vector<float> selfHostData = {0, 1, 2, 3}; std::vector<float> outHostData = {0, 0, 0, 0}; ret = CreateAclTensor(selfHostData, selfShape, &selfDeviceAddr, aclDataType::ACL_FLOAT, &self); CHECK_RET(ret == ACL_SUCCESS, return ret); ret = CreateAclTensor(outHostData, outShape, &outDeviceAddr, aclDataType::ACL_FLOAT, &out); CHECK_RET(ret == ACL_SUCCESS, return ret); // 3. 调用CANN算子库API（两段式） uint64_t workspaceSize = 0; aclOpExecutor* executor; // 第一段：获取workspace大小并完成入参校验 ret = aclnnSigmoidGetWorkspaceSize(self, out, &workspaceSize, &executor); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSigmoidGetWorkspaceSize failed. ERROR: %d\n", ret); return ret); // 根据第一段接口计算出的workspaceSize申请device内存 void* workspaceAddr = nullptr; if (workspaceSize > 0) { ret = aclrtMalloc(&workspaceAddr, workspaceSize, ACL_MEM_MALLOC_HUGE_FIRST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("allocate workspace failed. ERROR: %d\n", ret); return ret); } // 第二段：在指定stream上执行计算 ret = aclnnSigmoid(workspaceAddr, workspaceSize, executor, stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclnnSigmoid failed. ERROR: %d\n", ret); return ret); // 4. （固定写法）同步等待任务执行结束 ret = aclrtSynchronizeStream(stream); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("aclrtSynchronizeStream failed. ERROR: %d\n", ret); return ret); // 5. 将device侧结果拷贝回host并打印 auto size = GetShapeSize(outShape); std::vector<float> resultData(size, 0); ret = aclrtMemcpy(resultData.data(), resultData.size() * sizeof(resultData[0]), outDeviceAddr, size * sizeof(resultData[0]), ACL_MEMCPY_DEVICE_TO_HOST); CHECK_RET(ret == ACL_SUCCESS, LOG_PRINT("copy result from device to host failed. ERROR: %d\n", ret); return ret); for (int64_t i = 0; i < size; i++) { LOG_PRINT("result[%" PRId64 "] is: %f\n", i, resultData[i]); } // 6. 释放资源 aclDestroyTensor(self); aclDestroyTensor(out); aclrtFree(selfDeviceAddr); aclrtFree(outDeviceAddr); if (workspaceSize > 0) { aclrtFree(workspaceAddr); } aclrtDestroyStream(stream); aclrtResetDevice(deviceId); aclFinalize(); return 0; }

以上示例输入{0, 1, 2, 3}，理论输出分别为{0.5, 0.731059, 0.880797, 0.952574}（保留 6 位小数），可用于快速校验算子结果正确性。

6.3 aclnnInplaceSigmoid 与普通版的差异点

in-place 版本只需申请一个输入 Tensorself，第一段调用acclnInplaceSigmoidGetWorkspaceSize(self, &workspaceSize, &executor)，执行完成后直接从self对应的 Device 内存（selfDeviceAddr）拷回结果即可，无需构造out，资源释放时也无需释放输出 Tensor：

// 调用第一段接口 ret = aclnnInplaceSigmoidGetWorkspaceSize(self, &inplaceWorkspaceSize, &inplaceExecutor); // 按需申请workspace并调用第二段接口 ret = aclnnInplaceSigmoid(inplaceWorkspaceAddr, inplaceWorkspaceSize, inplaceExecutor, stream); // 同步后，直接从selfDeviceAddr读取结果 ret = aclrtMemcpy(inplaceResultData.data(), ..., selfDeviceAddr, ..., ACL_MEMCPY_DEVICE_TO_HOST);

七、源码级实现剖析：从 aclnn 到 Kernel 的完整链路

7.1 接口层（op_api）：类型适配与图编排

op_api/aclnn_sigmoid.cpp 是 aclnn 接口的核心实现，其执行编排逻辑集中在ExecSigmoidGetWorkspaceSize（L126-L172），主要步骤为：

1. 创建 OpExecutor：CREATE_EXECUTOR()固定写法；
2. 参数校验：CheckParams/CheckInplaceParams，对应上文 5.5 节的错误码；
3. 空 Tensor 短路：当self->IsEmpty()时直接返回 0 workspace，由 Kernel 侧处理空 Tensor（文档标注“支持空 Tensor”）；
4. 连续性处理：l0op::Contiguous(self, ...)将非连续输入转换为连续 Tensor；
5. 类型适配：若输入类型不在 Kernel 输出支持列表中（如 INT8/BOOL 等），则用l0op::Cast将输入转换为目标计算类型；
6. Kernel 调用：l0op::Sigmoid(selfCast, ...)执行真正的 Sigmoid 计算；
7. 结果回写：l0op::Cast将结果转换回out的数据类型，再通过l0op::ViewCopy拷贝到out（兼容非连续输出 Tensor）。

由此可见，接口层通过 Cast + Contiguous + ViewCopy 的组合，让 aclnn 接口能覆盖远比 Kernel 本身更广的数据类型与内存布局。

7.2 算子分派层（op_api/sigmoid.cpp）：AI Core 与 AI CPU 双路径

op_api/sigmoid.cpp 中定义了底层算子l0op::Sigmoid，其分派策略非常清晰：

• AICORE_DTYPE_SUPPORT_LIST = {DT_FLOAT, DT_FLOAT16, DT_BF16}：当输入类型为 FLOAT/FLOAT16/BF16 时，走AI Core 路径（SigmoidAiCore，通过ADD_TO_LAUNCHER_LIST_AICORE宏把算子加入 AI Core 任务队列）；
• 其余类型（DOUBLE、COMPLEX 等）走AI CPU 路径（SigmoidAiCpu，通过ADD_TO_LAUNCHER_LIST_AICPU宏下发）。

输出 Tensor 由executor->AllocTensor(self->GetViewShape(), self->GetDataType(), FORMAT_ND)按输入 shape 与类型自动分配，因此输出 shape 与输入天然一致。

7.3 Host 侧（op_host）：算子注册、shape 推导与 Tiling

• 算子定义：op_host/sigmoid_def.cpp 通过OpDef声明输入x、输出y均支持DT_FLOAT16/DT_BF16/DT_FLOAT与ND格式，并注册 ascend950、ascend350 两个平台的 AI Core 配置，同时开启DynamicShapeSupportFlag(true)（动态 shape 支持）、PrecisionReduceFlag(true)等特性，指定 Kernel 文件为sigmoid_apt。
• shape 推导：op_host/sigmoid_infershape.cpp 直接复用InferShape4Elewise逐元素通用推导，保证y.shape == x.shape。
• Tiling 计算：op_host/arch35/sigmoid_tiling_arch35.cpp 负责在 Host 侧根据输入 dtype 生成 TilingKey 与分块参数：dtype 组合（F16/F16、BF16/BF16、F32/F32）分别映射到OP_KEY_1/2/3（L54-L70），并通过GenerateTilingKey组合成最终 TilingKey；同时通过PlatformAscendC获取 AI Vector Core 数量（GetCoreNumAiv）与 UB 内存大小（GetCoreMemSize(CoreMemType::UB)），据此计算每个核分块与 UB 内循环次数。

7.4 Kernel 侧（op_kernel）：基于 TilingKey 的多态分派

op_kernel/sigmoid_apt.cpp 是 AI Core 侧的 Kernel 入口（__global__ __aicore__ void sigmoid），通过三个宏定义的 TilingKey 在设备侧进行多态分派：

#define SIGMOID_F16_TILING_KEY 100000000000100 // x是float16，y是float16 #define SIGMOID_BF16_TILING_KEY 200000000000100 // x是bfloat16，y是bfloat16 #define SIGMOID_F32_TILING_KEY 300000000000100 // x是float32，y是float32

Kernel 主体仅执行 AI Vector Core（KERNEL_TASK_TYPE_DEFAULT(KERNEL_TYPE_AIV_ONLY)），根据 TilingKey 分别实例化SigmoidF16、SigmoidBf16、SigmoidF32三种计算算子。三种 dtype 的具体实现位于 op_kernel/arch35/ 目录（sigmoid_f16.h / sigmoid_bf16.h / sigmoid_f32.h）。

以 op_kernel/arch35/sigmoid_f32.h 为例，可以看到典型的 AscendC 向量算子流水设计：

• Init阶段：设置 GlobalTensor 的全局内存地址（SetGlobalBuffer），并按elemNum初始化输入/输出双缓冲队列（DOUBLE_BUFFER = 2，即乒乓缓冲隐藏搬运延迟）；
• Process阶段：按块号（GetBlockIdx）判断当前是否为最后一个核，从而选取ubLoopOfFormerBlock/ubLoopOfTailBlock与ubTailOfFormerBlock/ubTailOfTailBlock做循环次数与尾块处理；每个循环依次执行CopyIn0（DataCopyPad从 GM 拷贝到 UB）→Compute1（向量计算）→CopyOut2（结果写回 GM）；
• 计算过程使用__VEC_SCOPE__向量作用域与寄存器级RegTensor/MaskReg优化，实现 Sigmoid 的向量化计算。

该设计使 Sigmoid 在 AI Core 上具备多核并行 + 双缓冲流水 + 尾块处理的完整性能优化结构，同时通过 Host 侧 Tiling 与设备侧 TilingKey 的配合支持动态 shape。

八、测试与验证

仓库为 Sigmoid 提供了完整的单元测试与系统测试体系，是验证算子正确性的最佳参照：

• 单元测试（ut）：位于 activation/sigmoid/tests/ut：
  • op_api/test_aclnn_sigmoid.cpp：aclnn 接口级测试；
  • op_host/test_sigmoid_infershape.cpp：shape 推导测试；
  • op_host/arch35/test_sigmoid_tiling.cpp：Tiling 计算测试；
  • op_kernel/test_sigmoid_apt.cpp：Kernel 层测试。
• 系统测试（st）：位于 activation/sigmoid/tests/st：
  • aclnnSigmoid/atk_aclnnSigmoid.json与executor_aclnnSigmoid.py：ATK（Ascend Test Kit）驱动的 aclnn 接口系统测试；
  • arch35/ttk_aclnn_sigmoid_st.csv、ttk_aclnn_inplace_sigmoid_st.csv、ttk_kernel_sigmoid_st.csv：TTK 驱动的接口与 Kernel 测试用例表；
  • assets/golden.py：Golden 数据生成脚本，用于按公式1 / (1 + exp(-x))生成预期结果。

测试数据从公式层面验证了算子输出与数学定义的一致性，同时覆盖了普通接口、in-place 接口与 Kernel 三条调用路径，可作为二次开发时回归测试的基线。

九、总结

CANN ops-nn 中的 Sigmoid 算子（activation/sigmoid）是一个“小而全”的典型逐元素算子：功能上严格遵循out = 1/(1+e^{-input})，产品上覆盖 Ascend 950 系列与 Atlas A2/A3/推理/训练全系列（Atlas 200I/500 A2 除外）；调用上同时支持aclnnSigmoid、acclnInplaceSigmoid两段式接口与图模式 IR 构图；实现上形成了aclnn 接口层（类型适配/连续性处理）→ l0op 分派层（AI Core / AI CPU 双路径）→ Host Tiling（多核分块）→ AI Core Kernel（TilingKey 多态分派 + 双缓冲流水）的完整链路。

对于开发者而言，理解本算子至少有三个层面的价值：一是掌握 aclnn 两段式接口的标准调用模板，可直接迁移到仓库内其他逐元素算子（如 gelu、relu 等）；二是理解“接口层宽类型 + Kernel 层窄类型 + Cast 桥接”的分层设计思想；三是借鉴其 TilingKey 分派与双缓冲 Kernel 写法，用于开发新的高性能 AI Core 算子。

【免费下载链接】ops-nn本项目是CANN提供的神经网络类计算算子库，实现网络在NPU上加速计算。项目地址: https://gitcode.com/cann/ops-nn