认证要求

实现Ascend C算子Sigmoid,算子命名为SigmoidCustom,编写其kernel侧代码、host侧代码,并完成aclnn算子调用测试。
相关算法:

$sigmoid(x) = 1/(1 + exp(-x))$

实现流程

环境初始配置

由于华为的ModelAtrs平台每次重启会把环境中除work目录之外的内容全部清除，因此每次需要重新设置环境，执行

bash init_env.sh
source ~/.bashrc
source /home/ma-user/Ascend/ascend-toolkit/set_env.sh
chmod +x -R *

上述命令依次用于下载包、配置环境变量以及脚本权限等

代码编写

从认证平台下载压缩包，解压之后包含SigmoidCustom和AclNNInvocation两个文件夹，SigmoidCustom中的内容主要完成代码编写、编译运行等，AclNNInvocation用于实现调用、测试等

文件主要内容如下

Host测实现

首先在sigmoid_custom_tiling.h完成tiling结构体设计，在代码的注释后添加如下内容

1 2	TILING_DATA_FIELD_DEF(uint32_t, totalLength); TILING_DATA_FIELD_DEF(uint32_t, tileNum);

然后在sigmoid_custom.cpp中的注释后添加

uint32_t totalLength = context->GetInputShape(0)->GetOriginShape().GetShapeSize();
context->SetBlockDim(BLOCK_DIM);
tiling.set_totalLength(totalLength);
tiling.set_tileNum(TILE_NUM);
tiling.SaveToBuffer(context->GetRawTilingData()->GetData(), context->GetRawTilingData()->GetCapacity());
context->GetRawTilingData()->SetDataSize(tiling.GetDataSize());
size_t *currentWorkspace = context->GetWorkspaceSizes(1);
currentWorkspace[0] = 0;
return ge::GRAPH_SUCCESS;

Device测实现

这里的代码稍微复杂，需要用户自行实现核函数的初始化、内存迁移、算法实现等，完整代码如下

#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;
class KernelSigmoid {
public:
    __aicore__ inline KernelSigmoid() {}
    __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t totalLength, uint32_t tileNum)
    {
        //考生补充初始化代码
        this->blockLength = totalLength / AscendC::GetBlockNum();
        this->tileNum = tileNum;
        this->tileLength = this->blockLength / tileNum / BUFFER_NUM;

        xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);
        yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + this->blockLength * AscendC::GetBlockIdx(), this->blockLength);

        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(half));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(half));
        pipe.InitBuffer(tmpBuffer1, this->tileLength * sizeof(half));
        pipe.InitBuffer(tmpBuffer2, this->tileLength * sizeof(half));
        pipe.InitBuffer(tmpBuffer3, this->tileLength * sizeof(half));
        pipe.InitBuffer(tmpBuffer4, this->tileLength * sizeof(half));
    }
    __aicore__ inline void Process()
    {
        //考生补充对“loopCount”的定义，注意对Tiling的处理
        int32_t loopCount = this->tileNum * BUFFER_NUM;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    __aicore__ inline void HighPrecisionReciprocal(LocalTensor<half>& dst, 
                                              const LocalTensor<half>& src, 
                                              int32_t length, 
                                              int iterations = 2)
    {
        LocalTensor<half> tmp = tmpBuffer4.Get<half>();
        half two = 2.0h,negone=-1.0h;

        // 初始近似值 (可以使用硬件Reciprocal的初始值)
        AscendC::Reciprocal(dst, src, length);

        // 牛顿迭代: x_{n+1} = x_n * (2 - a * x_n)
        for (int i = 0; i < iterations; ++i) {
            AscendC::Mul(tmp, src, dst, length);    // tmp = a * x_n
            AscendC::Muls(tmp, tmp, negone, length);   // tmp = 2 - a * x_n
            AscendC::Adds(tmp, tmp, two, length);
            AscendC::Mul(dst, dst, tmp, length);     // x_{n+1} = x_n * tmp
        }
    }

    __aicore__ inline void CopyIn(int32_t progress)
    {
        //考生补充算子代码
        AscendC::LocalTensor<half> xLocal = inQueueX.AllocTensor<half>();
        AscendC::DataCopy(xLocal, xGm[progress * this->tileLength], this->tileLength);
        inQueueX.EnQue(xLocal);
    }
    __aicore__ inline void Compute(int32_t progress)
    {
        //考生补充算子计算代码
        LocalTensor<half> xLocal = inQueueX.DeQue<half>();
        LocalTensor<half> yLocal = outQueueY.AllocTensor<half>();
        
        LocalTensor<half> tmp1 = tmpBuffer1.Get<half>();
        LocalTensor<half> tmp2 = tmpBuffer2.Get<half>();
        LocalTensor<half> tmp3 = tmpBuffer3.Get<half>();
        
        half one=1.0,negone=-1.0;
        AscendC::Muls(tmp1, xLocal, negone, this->tileLength);
        AscendC::Exp(tmp2, tmp1, this->tileLength);
        AscendC::Adds(tmp3, tmp2, one, this->tileLength);
        
        HighPrecisionReciprocal(yLocal,tmp3,this->tileLength,2);

        outQueueY.EnQue<half>(yLocal);
        inQueueX.FreeTensor(xLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress)
    {
        //考生补充算子代码
        AscendC::LocalTensor<half> yLocal = outQueueY.DeQue<half>();
        AscendC::DataCopy(yGm[progress * this->tileLength], yLocal, this->tileLength);
        outQueueY.FreeTensor(yLocal);
    }

private:
    TPipe pipe;
    //create queue for input, in this case depth is equal to buffer num
    TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
    //create queue for output, in this case depth is equal to buffer num
    TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
    GlobalTensor<half> xGm;
    GlobalTensor<half> yGm;

    //考生补充自定义成员变量
    TBuf<QuePosition::VECCALC> tmpBuffer1, tmpBuffer2, tmpBuffer3, tmpBuffer4;
    uint32_t blockLength;
    uint32_t tileNum;
    uint32_t tileLength;

};
extern "C" __global__ __aicore__ void sigmoid_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    KernelSigmoid op;
    //补充init和process函数调用内容
    op.Init(x, y, tiling_data.totalLength, tiling_data.tileNum);
    op.Process();
}

代码解释

一个自定义算子的生命周期

在深入代码之前，我们先理解一下从编写到执行，这个 SigmoidCustom 算子经历了什么：

开发阶段: 您编写了算子定义、Tiling 逻辑和 Kernel 实现这几个文件。
编译阶段: 使用 CANN 提供的工具链进行编译。编译器会分别处理 Host 端代码（编译成标准的 CPU 可执行指令）和 Device 端代码（编译成 AI Core 能理解的 CUBE 指令）。同时，OP_ADD(SigmoidCustom) 宏会将算子的元信息（如输入输出、Tiling 函数地址等）注册到昇腾的算子库中。
图构建与下发阶段: 当用户在一个神经网络模型（例如用 MindSpore 或 PyTorch）中调用 SigmoidCustom 时，前端框架会将模型转换为昇腾的计算图（Graph）。图编译器（Graph Engine, GE）会解析这个图，当遇到 SigmoidCustom 节点时，它会：

a. 调用 InferShape 和 InferDataType 函数，在不运行实际计算的情况下，静态推断出这个节点输出的形状和数据类型，以便为整个网络预先分配好内存。

b. 调用 TilingFunc 函数，根据当前节点的输入张量大小，生成具体的任务切分方案（即 TilingData）。

c. 将编译好的 Kernel CUBE 指令和 TilingData 打包成一个任务，下发给 NPU 的任务调度器（Task Scheduler）。
执行阶段: NPU 的任务调度器根据 TilingData 中的 blockDim 信息，唤醒指定数量的 AI Core。每个被唤醒的 AI Core 都会接收到 Kernel CUBE 指令和 TilingData 的一份拷贝，然后从 sigmoid_custom 这个 __global__ 入口函数开始，独立地执行计算任务。

Host 端代码深度解析 (CPU 侧)

Host 端代码的核心目标是**“定义规则”和“制定计划”**，它不参与实际的数值计算。

1.1 算子接口定义

这部分代码是 SigmoidCustom 算子与 CANN 框架交互的“接口”或“驱动”。框架本身不知道你这个算子具体怎么算，但通过这个定义，它知道了如何使用你。

this->Input("x")... 和 this->Output("y")...: 这不仅仅是定义了输入输出的名称，更重要的是定义了数据依赖关系。图编译器根据这个关系来构建计算图的拓扑结构，决定算子执行的先后顺序。
.DataType({ge::DT_FLOAT16}).Format({ge::FORMAT_ND}): 这是算子的能力声明。它告诉框架：“我能处理 FP16 类型的数据，并且不关心输入的具体维度（ND 格式）。” 如果上游算子的输出格式不匹配，图编译器会自动插入一个转置（Transpose）或类型转换（Cast）算子来适配。
this->SetInferShape(ge::InferShape): 这是一个至关重要的静态分析钩子。在模型编译时，真实的输入数据还不存在，但其形状（Shape）是已知的。框架会调用 InferShape 函数，让算子根据输入的 Shape "预言"自己输出的 Shape。对于 Sigmoid 这种 Element-wise（逐元素）操作，输出形状与输入形状完全相同，所以实现很简单：*y_shape = *x1_shape;。这使得框架可以在运行前就规划好整个模型的内存布局。
this->AICore().SetTiling(optiling::TilingFunc): 这是将**“计划制定者”**（Tiling 函数）与算子本身绑定的关键。框架知道，每当需要执行这个算子时，都应该去调用 TilingFunc 来获取详细的执行计划。

1.2 Tiling 函数

Tiling 函数是 Host 端的智慧所在，它在 CPU 上运行，为即将开始的 Device 端计算制定精密的作战计划。

gert::TilingContext\* context: context 对象是框架传递给 Tiling 函数的**“情报手册”**。它包含了关于当前任务的所有上下文信息，最主要的就是 context->GetInputShape(0)，它让 Tiling 函数知道了待处理的数据有多大。
核心决策:
1. 总览全局: uint32_t totalLength = ...GetShapeSize() 获取了总工作量。
2. 分配兵力: context->SetBlockDim(BLOCK_DIM) 决定了要动用多少个 AI Core（这里硬编码为8）。这是一个关键的并行策略决策。如果 totalLength 很小，可能只用一个 AI Core 更高效（避免多核通信开销）；如果 totalLength 巨大，则可以用满所有可用的 Core。
3. 制定指令:
  - tiling.set_totalLength(totalLength): 把全局信息写入“指令书” (TilingData)。
  - tiling.set_tileNum(TILE_NUM): 决定每个 AI Core 内部的流水线深度或切分粒度。这个值会影响性能，需要根据计算量和数据搬运量的比例（计算访存比）来权衡。TILE_NUM 越大，tileLength 就越小，数据搬运更频繁，但每个计算任务的延迟也更低，有利于流水线隐藏延迟。
信息传递: tiling.SaveToBuffer(...) 将填满信息的 TilingData 结构体序列化为一段连续的内存块。这段内存块将和编译好的 Kernel 代码一起被发送到 Device 端，成为每个 AI Core 执行任务的依据。

Device 端代码深度解析 (AI Core 侧)

Device 端代码是**“一线士兵”**，它们根据 Host 发来的计划，在自己的阵地（AI Core）上高效地执行计算。

2.1 内存模型与数据并行 (SPMD)

GM_ADDR x, GM_ADDR y: GM_ADDR 是 Global Memory Address 的缩写，代表全局内存中的一个地址。x 和 y 是输入和输出张量在全局内存中的起始地址。
SPMD (Single Program, Multiple Data): 这是 AI Core 的核心工作模式。所有被唤醒的 AI Core 都执行同一份程序代码（Single Program），但处理的数据各不相同（Multiple Data）。
AscendC::GetBlockIdx(): 这是实现 SPMD 的关键。该函数返回当前 AI Core 的唯一ID（从0到 blockDim-1）。
xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + this->blockLength * AscendC::GetBlockIdx(), ...): 这行代码是数据划分的精髓。
- (__gm__ DTYPE_X *)x: 将全局内存地址 x 转换为一个指向 FP16 类型的指针。
- ... + this->blockLength * AscendC::GetBlockIdx(): 计算偏移量。
  - 对于 ID=0 的 Core，偏移量是 blockLength * 0 = 0，它处理从 x 开始的数据。
  - 对于 ID=1 的 Core，偏移量是 blockLength * 1，它处理紧接着第一块的数据。
  - …
- 通过这种方式，totalLength 的数据被完美地分割成了 blockDim 份，每个 AI Core 都只关心自己负责的那一段，实现了无冲突的数据并行。

2.2 本地内存与流水线 (`TPipe`, `TQue`)

AI Core 的性能源泉在于其内部极速的本地内存（LM）和向量计算单元。流水线的目标就是让计算单元永远不要因为等待数据而停工。

TPipe, TQue: 这不是硬件，而是 Ascend C 提供的软件抽象，用于简化对硬件 DMA（数据搬运单元）和本地内存的管理。
InitBuffer: 在 Init 方法中，pipe.InitBuffer(inQueueX, BUFFER_NUM, ...) 相当于在本地内存中划分出 BUFFER_NUM (即2) 块独立的区域，并由 inQueueX 这个队列对象来管理。
双缓冲 (Double Buffering) 详解: BUFFER_NUM=2 构成了双缓冲。想象有两块缓冲区 A 和 B：
1. CopyIn: 调用 inQueueX.AllocTensor<half>() 会从空闲的缓冲区队列中取出一个（比如A），然后 DataCopy 指令启动 DMA 将 GM 数据拷贝到 A。完成后 inQueueX.EnQue(A) 将 A 标记为“已填充，可供计算”。
2. Compute: 调用 inQueueX.DeQue<half>() 会从“已填充”队列中取出一个（比如A），然后计算单元对 A 中的数据进行计算。
3. 重叠执行: 当计算单元正在处理 A 的数据时，CopyIn 阶段可以同时向另一个空闲的缓冲区 B 中拷贝下一批数据。当 A 计算完成，B 也正好拷贝完成，计算单元可以马不停蹄地开始处理 B，而 DMA 则开始将 A 的计算结果写回 GM。
for 循环的流水线建立过程:
- 迭代 i=0: CopyIn(0) 开始（数据块0从GM->LM），Compute(0) 和 CopyOut(0) 等待。
- 迭代 i=1: CopyIn(1) 开始（数据块1从GM->LM），同时 Compute(0) 开始（处理数据块0）。CopyOut(1) 等待。
- 迭代 i=2: CopyIn(2) 开始，Compute(1) 开始，CopyOut(0) 开始（将数据块0的结果从LM->GM）。至此，三级流水线完全建立，DMA的读、计算单元、DMA的写三个硬件部件都在满负荷工作。

2.3 向量计算与硬件指令

AscendC::Muls, AscendC::Exp, AscendC::Adds: 这些函数被称为内置函数（Intrinsics）。它们在编译时会被直接映射为 AI Core 的一条或多条向量指令。
例如，AscendC::Exp(tmp2, tmp1, this->tileLength) 不是一个循环，而是一条指令。它告诉向量计算单元：“对 tmp1 指向的 tileLength 个半精度浮点数，逐个执行指数运算，并将结果存放到 tmp2 中”。这种 SIMD（单指令多数据） 的处理方式是 AI Core 性能的根本来源。

2.4 数值计算优化

问题: 硬件通常提供速度快但精度有限的近似计算指令，比如求倒数。对于某些需要高精度的算法，这可能导致误差累积。
解决方案: 牛顿-拉弗森迭代法是一种经典的数值分析方法，可以用来求方程的根。对于求 1/a，相当于求方程 f(x) = 1/x - a 的根。
代码映射:
1. AscendC::Reciprocal(dst, src, length);: 使用硬件指令得到一个初始近似值 x_0 (即 dst)。src 是 a。
2. AscendC::Mul(tmp, src, dst, length);: 计算 a * x_n，结果存入 tmp。
3. AscendC::Muls(tmp, tmp, negone, length);: 变为 - (a * x_n)。
4. AscendC::Adds(tmp, tmp, two, length);: 变为 2 - a * x_n。
5. AscendC::Mul(dst, dst, tmp, length);: 计算 x_n * (2 - a * x_n)，得到 x_{n+1}，并覆盖回 dst。
这个 for 循环每迭代一次，结果的精度就会翻倍。迭代2次就能获得非常高的精度，这体现了算法与硬件结合的优化思想。

运行测试

在运行之前，首先检查CMakePresets.json文件中的ASCEND_COMPUTE_UNIT和ASCEND_CANN_PACKAGE_PATH是否正确，如下

"ASCEND_COMPUTE_UNIT": {
    "type": "STRING",
    "value": "ascend310b;ascend910b"
},
"ASCEND_CANN_PACKAGE_PATH": {
    "type": "PATH",
    "value": "/home/ma-user/Ascend/ascend-toolkit/latest"
},

然后检查SigmoidCustom中的AICore是否设置正确，如下

1 2	this->AICore().AddConfig("ascend910b") .AddConfig("ascend310b");

如果上述设置没有问题，就可以进行编译测试了

运行编译

1	bash build.sh

编译之后会得到build_out目录，运行其中的安装包

1 2	cd build_out ./custom_opp_ubuntu_aarch64.run

最后到AclNNInvocation中运行测试

1	bash run.sh

检验脚本会调用py脚本，分别用numpy和我们写的算子计算答案，然后对比误差，如果代码没有问题会给出测试成功的提示。由于华为云提供的这个环境属实有点贵，因此我就不再开机测试并截图了，大家自行操作

随后按照认证要求将SigmoidCustom打包提交即可，测试通过后会获得认证证书，如下