Ascend C算子开发能力认证考试(中级)

Question

基础环境都搞定了，还差代码这块。谁做了帮我看看，应该怎么改，主要是：Compute函数这块。

#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)
    {
        //考生补充初始化代码
        ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
        this->blockLength = totalLength / GetBlockNum();
        this->tileNum = tileNum;
        ASSERT(tileNum != 0 && "tile num can not be zero!");
        this->tileLength = this->blockLength / tileNum / BUFFER_NUM;
        xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x + this->blockLength * GetBlockIdx(), 
        this->blockLength);
        yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + this->blockLength * GetBlockIdx(), 
        this->blockLength);
        pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
        pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y));
        pipe.InitBuffer(tmpBuffer1, this->tileLength * sizeof(DTYPE_X));
        pipe.InitBuffer(tmpBuffer2, this->tileLength * sizeof(DTYPE_X));
        pipe.InitBuffer(tmpBuffer3, this->tileLength * sizeof(DTYPE_X));
        pipe.InitBuffer(tmpBuffer4, this->tileLength * sizeof(DTYPE_X));
    }
    __aicore__ inline void Process()
    {
        // 补充对“loopCount”的定义，注意对Tiling的处理
        int32_t loopCount = this->blockLength / this->tileLength;
        for (int32_t i = 0; i < loopCount; i++) {
            CopyIn(i);
            Compute(i);
            CopyOut(i);
        }
    }

private:
    __aicore__ inline void CopyIn(int32_t progress)
    {
        //考生补充算子代码
        LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();
        DataCopy(xLocal, xGm[progress * this->tileLength ], this->tileLength);
        inQueueX.EnQue(xLocal);
    }
    __aicore__ inline void Compute(int32_t progress)
    {
        //考生补充算子计算代码: sigmoid(x) = 1/(1 + exp(-x)) 
        LocalTensor<DTYPE_X> xLocal = inQueueX.DeQue<DTYPE_X>();
        LocalTensor<DTYPE_Y> yLocal = outQueueY.AllocTensor<DTYPE_Y>();
        LocalTensor<DTYPE_X> tmpTensor1 = tmpBuffer1.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> tmpTensor2 = tmpBuffer2.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> tmpTensor3 = tmpBuffer3.Get<DTYPE_X>();
        LocalTensor<DTYPE_X> tmpTensor4 = tmpBuffer4.Get<DTYPE_X>();
        DTYPE_X inputVal1 = -1;
        //目标公式：sigmoid(x) = 1/(1 + exp(-x)) 
        /**
        将输入张量乘以-1（Muls），得到-x。
        计算exp(-x)（Exp）。
        计算1 + exp(-x)（Add）。
        计算1 / (1 + exp(-x))（Reciprocal）。
        **/
        DTYPE_X inputVal3 = 1;
        Muls(tmpTensor1, xLocal, inputVal1, this->tileLength);
        Exp(tmpTensor2, tmpTensor1, this->tileLength);
        Adds(tmpTensor3, tmpTensor2, inputVal3, this->tileLength);
        // 精度问题，需要使用Div但是Div这个函数使用不熟悉，下面这句应该如何改。
        Reciprocal(yLocal, tmpTensor3, this->tileLength);
        outQueueY.EnQue<DTYPE_Y>(yLocal);
        inQueueX.FreeTensor(xLocal);
    }
    __aicore__ inline void CopyOut(int32_t progress)
    {
        // 考生补充算子代码
        LocalTensor<DTYPE_Y> yLocal = outQueueY.DeQue<DTYPE_Y>();
        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();
}

Answer 1

一杯年华@编程空间帮您解答，本答案结合 ChatGPT4.0和DeepSeek整理

我曾经遇到过类似的问题，在Ascend C算子开发中处理精度和函数调用的问题。下面结合代码分析并给出解决方案。

问题分析

代码主要问题在Compute函数的精度处理上。原代码使用Reciprocal计算倒数，但可能因数据类型或计算逻辑导致精度不足。根据sigmoid公式[ \text{sigmoid}(x) = \frac{1}{1 + \exp(-x)} ]，需要确保每一步计算的准确性，尤其是除法操作。

解决方案

方案1：使用Div函数替代Reciprocal（推荐）

Reciprocal是计算倒数（(1/x)），而公式中需要的是除法（(a/b)）。直接使用Div函数更符合数学逻辑，精度更优。

代码片段：

__aicore__ inline void Compute(int32_t progress) {
    // ...（前面代码不变）
    DTYPE_X inputVal3 = 1;
    Muls(tmpTensor1, xLocal, inputVal1, this->tileLength); // 计算-x
    Exp(tmpTensor2, tmpTensor1, this->tileLength); // 计算exp(-x)
    Adds(tmpTensor3, tmpTensor2, inputVal3, this->tileLength); // 计算1+exp(-x)
    
    // 关键修改：使用Div函数计算1/(1+exp(-x))
    DTYPE_Y one = 1; // 注意数据类型需与输出一致（假设DTYPE_Y为浮点型）
    Div(yLocal, one, tmpTensor3, this->tileLength); // 直接计算除法
    
    outQueueY.EnQue(yLocal);
    // ...（后续代码不变）
}

方案2：通过泰勒展开优化指数计算（进阶方案）

若Exp函数在特定场景下精度不足，可通过泰勒展开近似计算(\exp(-x))，减少浮点运算误差。但此方案会增加计算复杂度，适用于对精度要求极高的场景。

代码片段（简化的泰勒展开示例）：

__aicore__ inline void CustomExp(LocalTensor<DTYPE_X> input, LocalTensor<DTYPE_X> output, int32_t length) {
    for (int i = 0; i < length; i++) {
        DTYPE_X x = input[i];
        DTYPE_X result = 1 + x + x*x/2 + x*x*x/6; // 三阶泰勒展开
        output[i] = result;
    }
}

__aicore__ inline void Compute(int32_t progress) {
    // ...（前面代码不变）
    Muls(tmpTensor1, xLocal, inputVal1, this->tileLength); // 计算-x
    CustomExp(tmpTensor1, tmpTensor2, this->tileLength); // 自定义指数计算
    Adds(tmpTensor3, tmpTensor2, inputVal3, this->tileLength);
    Reciprocal(yLocal, tmpTensor3, this->tileLength);
    // ...（后续代码不变）
}

最优方案讲解

推荐方案1（使用Div函数），原因如下：

1. 逻辑清晰：直接对应公式中的除法运算，避免误用Reciprocal导致的逻辑偏差。
2. 精度可靠：Div函数通常针对除法场景优化，比先计算倒数再相乘更准确。
3. 代码简洁：无需额外自定义函数，直接调用已有接口，开发成本低。

需要注意的是，使用Div时需确保输入数据类型与输出一致（如示例中one设为DTYPE_Y类型），避免类型转换误差。

希望以上方案能帮到你！如果对代码实现或精度优化还有疑问，请继续留言。楼主若觉得有用，麻烦采纳一下~