首页 > 其他分享 >Ascend C算子开发（中级）—— 编写Sinh算子

Ascend C算子开发（中级）—— 编写Sinh算子

时间：2024-09-08 22:52:20浏览次数：8  
标签：__ tileLength tileNum SinhCustom Ascend Sinh 算子 tiling

在这里插入图片描述
在这里插入图片描述

Ascend C算子开发（中级）—— 编写Sinh算子

文章目录

准备工作

一块香橙派AI pro开发板，一根Type-c口的电源线，一根网线，一个网线转接器，一台笔记本电脑。
在这里插入图片描述

香橙派与PC连接

1). 硬件连接与启动（如下图所示）

a）检查Orange Pi AI pro 是否已经插入SD卡；
b）使用网线连接Orange Pi AI pro以太网口，网线另一端连接PC/转接头；
c）连接电源，如下图所示两个LED指示灯绿色常亮，表示启动正常；
d）网口下方两个灯，右侧绿灯常亮，左侧橙灯闪烁，代表网口物理连接正常

2).从PC远程登录到香橙派（根据各自系统版本选择）

  • （Windows）以太网口远程登录：
    开发文档（备注，OPi AIpro以太网口为192.168.137.100，Windows PC以太网可设置为192.168.137.101 ； ssh链接以 root 用户名登录，密码为 Mind@123）

  • Mac系统远程登录：
    开发文档（备注，OPi AIpro以太网口为192.168.137.100，Mac PC以太网口可设置为192.168.137.2 ；ssh链接以 root 用户名登录，密码为 Mind@123）

Add 算子调用体验

编译并调用一个Add算子的全过程：

  • 编译
cd ~/samples/operator/AddCustomSample/FrameworkLaunch/AddCustom
bash build.sh

在这里插入图片描述

  • 部署
cd build_out
./custom_opp_ubuntu_aarch64.run

在这里插入图片描述

  • 调用
cd ~/samples/operator/AddCustomSample/FrameworkLaunch/AclNNInvocation
bash run.sh

在这里插入图片描述

Sinh算子开发（Ascend C算子开发中级认证考试内容）

1). 使用提供的考试代码工程，cd /root/SinhCustom/SinhCustom ，依次打开下图红框所示的三个源码文件，并根据注释提示补全代码；可以对照Add算子的例子进行修改。
在这里插入图片描述

  • op_host端

sinh_custom.cpp

#include "sinh_custom_tiling.h"
#include "register/op_def_registry.h"

namespace optiling {
/**
Tiling Func负责对输入数据进行分块（Tile）处理。分块处理的好处在于，可以并行计算不同块中的数据，提升计算效率。
BLOCK_DIM 定义了每次计算操作需要处理的块的数量。
TILE_NUM 定义了在每个计算块中进一步将数据划分为更小的子块。每个子块的数据大小由blocklength/TILE_NUM来决定。
该方法将 totalLength 和 TILE_NUM 此类方法保存在tiling对象中，随后将这些信息写入`RawTilingData`中
**/
static ge::graphStatus TilingFunc(gert::TilingContext* context)
{
    SinhCustomTilingData tiling;
    //考生自行填充
    const uint32_t BLOCK_DIM = 8;
    const uint32_t TILE_NUM = 8;
    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;
}
}
/**
这个函数定义了输入与输出的形状推理逻辑，保证输入和输出的形状是相同的。
**/
namespace ge {
static ge::graphStatus InferShape(gert::InferShapeContext* context)
{
    const gert::Shape* x1_shape = context->GetInputShape(0);
    gert::Shape* y_shape = context->GetOutputShape(0);
    *y_shape = *x1_shape;
    return GRAPH_SUCCESS;
}
}
/**
该类定义了一个自定义的sinh算子，明确了输入和输出的张量格式和数据类型（DT_FLOAT16），并且指定该算子的推理形状函数是InferShape，Tiling函数是TilingFunc。
最后，通过OP_ADD(SinhCustom)将该算子注册到Ascend编译器中。
**/
namespace ops {
class SinhCustom : public OpDef {
public:
    explicit SinhCustom(const char* name) : OpDef(name)
    {
        this->Input("x")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND});
        this->Output("y")
            .ParamType(REQUIRED)
            .DataType({ge::DT_FLOAT16})
            .Format({ge::FORMAT_ND})
            .UnknownShapeFormat({ge::FORMAT_ND});

        this->SetInferShape(ge::InferShape);

        this->AICore()
            .SetTiling(optiling::TilingFunc);
        this->AICore().AddConfig("ascend310b");
    }
};

OP_ADD(SinhCustom);
}

sinh_custom_tilling.h


#include "register/tilingdata_base.h"
/**
这里定义了tiling数据结构的字段totalLength和tileNum，它们分别表示输入数据的总长度和分块数目。通过REGISTER_TILING_DATA_CLASS将SinhCustomTilingData与算子SinhCustom进行绑定。
**/
namespace optiling {
BEGIN_TILING_DATA_DEF(SinhCustomTilingData)
  //考生自行定义tiling结构体成员变量
  TILING_DATA_FIELD_DEF(uint32_t, totalLength);
  TILING_DATA_FIELD_DEF(uint32_t, tileNum);
END_TILING_DATA_DEF;
REGISTER_TILING_DATA_CLASS(SinhCustom, SinhCustomTilingData)
}
  • op_kernel端

前两个类和Add的算子对应类完全相同，关键需要修改的是op_kernel端的逻辑，因为sinh算子的公式为sinh(x) = (exp(x) - exp(-x)) / 2.0，总共分为四个部分，分别是

sinh_custom.cpp

#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;

class KernelSinh {
public:
    __aicore__ inline KernelSinh() {}
    /**
    该函数负责初始化全局和局部缓存、块和Tile的长度，并根据tileNum和blockLength来计算tileLength。
xGm.SetGlobalBuffer 和 yGm.SetGlobalBuffer 初始化全局内存上的输入和输出数据区域。
pipe.InitBuffer 初始化了多个队列和临时缓冲区，用于算子执行过程中数据的缓存和处理。
    **/
    __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()
    {
        /*
        Process函数执行主循环，每次循环中执行三个步骤：从全局内存拷贝数据到局部内存（CopyIn），计算（Compute），然后将结果从局部内存拷贝回全局内存（CopyOut）。
        */
        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 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)
    {
        //考生补充算子计算代码
        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;
        DTYPE_X inputVal2 = 0.5;
        //sinh(x) = (exp(x) - exp(-x)) / 2.0
        /**
        将输入张量乘以-1（Muls），得到-x。
		计算exp(-x)（Exp）。
		计算exp(x)。
		计算exp(x) - exp(-x)（Sub）。
		将结果乘以0.5，得到sinh(x)的结果（Muls）。
        **/
        Muls(tmpTensor1, xLocal, inputVal1, this->tileLength);
        Exp(tmpTensor2, tmpTensor1, this->tileLength);
        Exp(tmpTensor3, xLocal, this->tileLength);
        Sub(tmpTensor4, tmpTensor3, tmpTensor2, this->tileLength);
        Muls(yLocal, tmpTensor4, inputVal2, 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;
};
/**
这是最终的自定义内核函数，通过Init函数初始化操作，并调用Process函数执行具体计算。
**/
extern "C" __global__ __aicore__ void sinh_custom(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    KernelSinh op;
    //补充init和process函数调用内容
    op.Init(x, y, tiling_data.totalLength, tiling_data.tileNum);
    op.Process();
}

2)代码补齐完成后，cd /root/SinhCustom/SinhCustom ，然后执行如下命令进行编译构造：
bash build.sh
当命令显示如下信息，证明构建成功
在这里插入图片描述
3)构建成功之后，

cd /root/SinhCustom/SinhCustom/build_out 
./custom_opp_ubuntu_aarch64.run

当命令行显示如下信息证明安装成功:

4)最后，

cd /root/SinhCustom/AclNNInvocation 
bash run.sh

当命令行显示如下信息，说明通过测试.

5)测试通过后，将上述代码打包在zip 包内，例如使用如下命令：

cd /root 
zip -r SinhCustom.zip SinhCustom

6)打包完成后，到MobaXTerm左侧的文件栏中找到压缩包，下载到PC本地，上传到考试页面，交卷即可。

在这里插入图片描述

标签：__,tileLength,tileNum,SinhCustom,Ascend,Sinh,算子,tiling
From： https://blog.csdn.net/m0_74120525/article/details/142027001

相关文章

  • 说说Canny边缘检测算子?
    Canny边缘检测算子什么是Canny边缘检测算子Canny边缘检测算子原理Canny边缘检测算子用途什么是Canny边缘检测算子Canny边缘检测算子是一种旨在以最优方式从图像中提取边缘信息的算法。其“最优”体现在三个方面：低错误率：算法应尽可能多地标识出图像中的实际边缘，同时......
  • 图像边缘检测技术详解：利用OpenCV实现Sobel算子
    图像边缘检测技术详解：利用OpenCV实现Sobel算子前言Sobel算子的原理代码演示结果展示结语前言  在数字图像处理的广阔领域中，边缘检测技术扮演着至关重要的角色。无论是在科学研究、工业自动化，还是在日常生活中的智能设备中，我们都需要从图像中提取有用的信息。边缘，作......
  • 【AI System】Ascend NPU 架构 & CANN 平台入门学习
    AscendNPU架构&CANN平台入门学习概述昇腾NPU是专门用于AI训练/推理计算的AI专用处理器，其中的AICore能够在很大程度上提高AI计算的效率。本文将主要介绍ASCENDNPU的硬件架构&工作原理、AICore的计算模式以及异构计算平台CANN等内容。NPU硬件架......
  • openGauss-资源池化算子卸载
    openGauss-资源池化算子卸载可获得性本特性自openGauss5.1.0版本开始引入。特性简介共享存储虽然带来弹性，可靠性的好处，但是和本地盘单机比较性能会下降较多，主要是网络IO和分布式存储自身带来的延迟，尤其对于大规模查询bufferpool无法缓存的场景，大量的数据需要从存储节点搬......
  • [ARC183E] Ascendant Descendant
    MyBlogs[ARC183E]AscendantDescendant一个直接的想法是求出\(L_i,R_i\)表示极大的区间\([L_i,R_i]\)满足\(\forallj\in[L_i,R_i],b_j\in\text{subtree}(a_i)\)。由于树的性质，\([L_i,R_i]\)之间要么相离，要么包含。但是\(L_i,R_i\)并不是\(i\)能真正到达的点。因......
  • 昇腾 - AscendCL C++应用开发 线程安全的队列
    昇腾-AscendCLC++应用开发线程安全的队列flyfishC++mutex各种各样的互斥锁mutex、timed_mutex、recursive_mutex、shared_mutexC++线程间同步的条件变量std::condition_variable和std::condition_variable_anyC++提供的智能指针unique_ptr、shared_ptr、wea......
  • 昇腾 - AscendCL C++应用开发 目标检测中的非极大值抑制NMS和计算候选边界框之间的交
    昇腾-AscendCLC++应用开发目标检测中的非极大值抑制（NMS，Non-MaximumSuppression）涉及计算候选边界框之间的交并比（IOU，IntersectionoverUnion）flyfish结构体BBox:定义了一个边界框的数据结构，包含中心坐标、宽高、置信度分数、类别索引和输出索引。函数IOU:计算两个......
  • 03-RDD算子
    1.转换类算子1.1基本转换类算子1、Map新的RDD中的每个元素与旧的RDD的每个元素是一对一的关系。objectCH_0201_RDDAPI_Map{defmain(args:Array[String]):Unit={valconf:SparkConf=newSparkConf().setMaster("local").setAppName("Map")valsc=ne......
  • 配置 昇腾 Ascend C/C++ 开发环境
    配置昇腾AscendC/C++开发环境flyfish这里以OrangePiAiPro为例先说如何配置MindStudio，然后再说如何查看OrangePiAiPro的一些信息OrangePiAIPro开发板是香橙派联合华为精心打造的高性能AI开发板，其搭载了昇腾AI处理器。Linux桌面系统的默认登录用户为H......
  • 深度解读昇腾CANN小shape算子计算优化技术，进一步减少调度开销
    摘要：Host调度模式下，GE将模型中算子的执行单元划分为HostCPU执行与Device（昇腾AI处理器）执行两大类。本文分享自华为云社区《深度解读昇腾CANN小shape算子计算优化技术，进一步减少调度开销》，作者：昇腾CANN。GE（GraphEngine）将模型的调度分为Host调度与下沉调度两种模式。经过上期的介......

赞助商

阅读排行

网站主页关于我们联系我们网站地图

本网站内容转载自其他媒体，侵权联系[admin##ips99.com]。

Copyright © 2020-2023 IPS99 版权所有 IPS99