[嵌入式AI从0开始到入土]17_Ascend C算子开发

[嵌入式AI从0开始到入土]嵌入式AI系列教程

注:等我摸完鱼再把链接补上
可以关注我的B站号工具人呵呵的个人空间,后期会考虑出视频教程,务必催更,以防我变身鸽王。

第1期 昇腾Altas 200 DK上手
第2期 下载昇腾案例并运行
第3期 官方模型适配工具使用
第4期 炼丹炉的搭建(基于Ubuntu23.04 Desktop)
第5期 炼丹炉的搭建(基于wsl2_Ubuntu22.04)
第6期 Ubuntu远程桌面配置
第7期 下载yolo源码及样例运行验证
第8期 在线Gpu环境训练(基于启智ai协作平台)
第9期 转化为昇腾支持的om离线模型
第10期 jupyter lab的使用
第11期 yolov5在昇腾上推理
第12期 yolov5在昇腾上应用
第13期_orangepi aipro开箱测评
第14期 orangepi_aipro小修补含yolov7多线程案例
第15期 orangepi_aipro欢迎界面、ATC bug修复、镜像导出备份
第16期 ffmpeg_ascend编译安装及性能测试
第17期 Ascend C算子开发
未完待续…


文章目录

  • [嵌入式AI从0开始到入土]嵌入式AI系列教程
  • 前言
  • 一、环境配置
    • 1、CANN包安装
    • 2、配置ssh密钥(可选)
    • 3、配置git(可选)
  • 二、获取sample样例
    • 1、add算子
      • 1、KernelLaunch
      • 2、Framework
      • 3、AclNN
    • 2、Addcdiv算子
  • 三、编写自己的算子
    • 1、搭建框架
    • 2、 KernelLaunch编写
      • 1、myCustom.cpp
      • 2、main.cpp
      • 3、scripts/gen_data.py
    • 3、 framework编写
    • 4、 Aclnn测试
  • 四、torch_npu重新编译(可选)
  • 五、常用api
  • 问题
    • 1、fatal error: register/tilingdata_base.h: No such file or directory
  • 总结


前言

我在24年3月和我的小伙伴一起参加了第一届昇腾AI原生创新精英挑战赛,在这里做一下总结。这里以orangepi Ai Pro为例。
注:我们的代码仓最早将于24.05.10开放,大家可以直接看op_kernel内的compute,kernelLaunch内可能有错,实在来不及改了
代码仓地址:https://gitee.com/toolsmanhehe/acl_ops

一、环境配置

我们基于正常能够使用的镜像作为基础镜像。这里我推荐使用minimal镜像。这样就不用先卸载cann了,甚至你可以直接删除/opt/compress目录,反正咱后面直接远程连接敲代码了,也用不上。

1、CANN包安装

wget https://ascend-repo.obs.cn-east-2.myhuaweicloud.com/Milan-ASL/Milan-ASL%20V100R001C17SPC702/Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run
chmod +x Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run

#卸载旧的CANN
./Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run --uninstall
sudo rm -rf /usr/local/Ascend/ascend-toolkit/*

#安装指定版本的CANN
./Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run --install
#安装依赖
pip install protobuf==3.20.0
#添加环境变量
echo “source /usr/local/Ascend/ascend-toolkit/set_env.sh” >> /home/HwHiAiUser/.bashrc
source /home/HwHiAiUser/.bashrc

2、配置ssh密钥(可选)

主要是vscode等ide连接时都需要输入密码,比较麻烦。
这里可以参考我之前的文章来实现免密登录,在七、问题 的第5点

3、配置git(可选)

因为我们三个人在三个城市,因此为了方便讨论和开发,我们建立了代码仓库,但是每次推送和拉取都需要账号密码(在完赛前是不可能公开的),这不符合本懒人的性格啊。
这里我们需要在开发环境上执行

cd ~
touch .git-credentials
vim .git-credentials
#输入以下内容,请自行替换username和password
https://username:password@gitee.com
git config --global credential.helper store

二、获取sample样例

cd 
git clone https://gitee.com/ascend/samples.git

在不修改算子名称,输入输出的时候,我们只需要关注图中框出来的文件即可。

1、add算子

打开目录operator/AddCustomSample

1、KernelLaunch

在这里插入图片描述
我们的调用顺序是main.cpp->add_custom_do->add_custom->op.Init->op.Process
在这里插入图片描述
因为我们要实现的算子的Z=X+Y,因此我们需要将这三个变量传入计算过程。
虽然这里只有2个输入,但是输出也需要申请内存,因此是3个输入参数

然后我们需要申明相关的变量和常量(这里使用静态shape)。
在这里插入图片描述
在这里插入图片描述
接着就是初始化,为各个张量申请内存
在这里插入图片描述
接着就是计算过程,这里因为使用的是静态shape,因此循环次数是定值(芯片内存空间有限,不可能一次性全部计算完成)
在这里插入图片描述
在copyin的时候从xGM和yGM分别取出TILE_LENGTH个数据,存入xLocal和yLocal以供compute使用。
在compute结束以后,我们需要先使用outQueueZ.EnQue来表示计算完成,但是此时不能释放zLocal的内存,因为我们还没有保存到zGM。
在copyout环节,将输出结果存入zGM。

接着我们看生成测试数据的程序,这里我们生成了2条16384个1~100随机half格式的数据。我们最后可以直接对比output/golden.binoutput/output_z.bin的md5值来判断算子正确与否。或者修改scripts/verify_result.py直接打印误差数量。
在这里插入图片描述
最后来到KernelLaunch目录执行以下命令,测试核函数正确性。
务必先进行cpu测试,通过后执行npu测试,在npu下有些报错不显示

su			   #使用root用户执行,否则可能报错
bash run.sh -r cpu -v ascend310B1   #cpu测试
bash run.sh -r npu -v ascend310B1   #npu测试

以下为cpu测试结果
在这里插入图片描述
以下为npu测试结果
在这里插入图片描述
测试均通过的情况下,我们就可以进行下一步的framework的编写了

2、Framework

我们先看AddCustomSample/FrameworkLaunch/AddCustom.json这个文件,上面为输入变量,下面为输出变量。我们需要使用这个配置文件来生成framework工程。此处的变量应该和工程内的一致。
在这里插入图片描述
接着我们看工程。
在这里插入图片描述
op_host没什么可说的,可以去看本文下一个案例Addcdiv。
op_kernel基本上就是把上面在kernelLaunch中测试通过的代码cv过来。
注意图中的地方就可以了,这个tiling是从host侧传入的。然后在开头将静态shape删除了,因为这里我们是通过op_host实现的动态shape的切分,然后传入kernel侧的。
在这里插入图片描述
接下来修改CMakePresets.json,将框出来的地方改成你的CANN路径。
在这里插入图片描述
最后,我们进入framework目录,编译算子并安装

bash build.sh
./build_out/custom_opp_ubuntu_aarch64.run

在这里插入图片描述

3、AclNN

在算子大赛的时候,这个是由官方发布的(就是可能有错误),我们直接使用即可,一般测试能通过,就会有4-8分(10分满分)。
在这里插入图片描述
这里的gen_data和kernelLaunch里是一样的,我们执行以下命令,验证算子正确与否。

bash run.sh

测试通过会有如下提示
在这里插入图片描述

2、Addcdiv算子

打开目录operator/AddcdivCustomSample
大部分与add算子相似,因此我们这里只看op_host和op_kernel部分。
在头文件中你会发现多了许多东西,所有的东西我们都需要传入kernel侧。具体实现过程就去阅读代码吧,就是这个案例也是赶出来的,可能里面的切分策略不是最好的,但是确实是能用的。

#ifndef ADDCDIV_CUSTOM_TILING_H
#define ADDCDIV_CUSTOM_TILING_H
#include "register/tilingdata_base.h"

namespace optiling {
BEGIN_TILING_DATA_DEF(AddcdivCustomTilingData)
  TILING_DATA_FIELD_DEF(float, value); 	//参与计算的标量
  TILING_DATA_FIELD_DEF(uint32_t, blockLength);
  TILING_DATA_FIELD_DEF(uint32_t, tileNum);
  TILING_DATA_FIELD_DEF(uint32_t, tileLength);
  TILING_DATA_FIELD_DEF(uint32_t, lasttileLength);
  TILING_DATA_FIELD_DEF(uint32_t, formerNum);
  TILING_DATA_FIELD_DEF(uint32_t, formerLength);
  TILING_DATA_FIELD_DEF(uint32_t, formertileNum);
  TILING_DATA_FIELD_DEF(uint32_t, formertileLength);
  TILING_DATA_FIELD_DEF(uint32_t, formerlasttileLength);
  TILING_DATA_FIELD_DEF(uint32_t, tailNum); 
  TILING_DATA_FIELD_DEF(uint32_t, tailLength);
  TILING_DATA_FIELD_DEF(uint32_t, tailtileNum);
  TILING_DATA_FIELD_DEF(uint32_t, tailtileLength);
  TILING_DATA_FIELD_DEF(uint32_t, taillasttileLength);    
END_TILING_DATA_DEF;

REGISTER_TILING_DATA_CLASS(AddcdivCustom, AddcdivCustomTilingData)
}
#endif // ADDCDIV_CUSTOM_TILING_H

以下为op_kernel内的部分代码

 private:
  TPipe pipe;
  // TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY, inQueueZ;
  TQue<QuePosition::VECIN, BUFFER_NUM> inQueueIN;
  TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOUT;
  GlobalTensor<half> xGm;
  GlobalTensor<half> yGm;
  GlobalTensor<half> zGm;
  GlobalTensor<half> outGm;
  half value;
  uint32_t blockLength;
  uint32_t tileNum;
  uint32_t tileLength;
  uint32_t lasttileLength;
  uint32_t formerNum;
  uint32_t formerLength;
  uint32_t formertileNum;
  uint32_t formertileLength;
  uint32_t formerlasttileLength;
  uint32_t tailNum;
  uint32_t tailLength;
  uint32_t tailtileNum;
  uint32_t tailtileLength;
  uint32_t taillasttileLength;
};

extern "C" __global__ __aicore__ void addcdiv_custom(GM_ADDR x, GM_ADDR y,
                                                     GM_ADDR z, GM_ADDR out,
                                                     GM_ADDR workspace,
                                                     GM_ADDR tiling) {
  GET_TILING_DATA(tiling_data, tiling);
  // TODO: user kernel impl
  KernelAddcdiv op;

  uint32_t tilingKey = 1;
  if (TILING_KEY_IS(1)) {
    tilingKey = 1;
  } else if (TILING_KEY_IS(2)) {
    tilingKey = 2;
  } else {
    tilingKey = 1;
  }

  op.Init(x, y, z, out, tiling_data.value, tiling_data.blockLength,
          tiling_data.tileNum, tiling_data.tileLength,
          tiling_data.lasttileLength, tiling_data.formerNum,
          tiling_data.formerLength, tiling_data.formertileNum,
          tiling_data.formertileLength, tiling_data.formerlasttileLength,
          tiling_data.tailNum, tiling_data.tailLength, tiling_data.tailtileNum,
          tiling_data.tailtileLength, tiling_data.taillasttileLength,
          tilingKey);
  op.Process();
}

#ifndef __CCE_KT_TEST__
// call of kernel function
void addcdiv_custom_do(uint32_t blockDim, void* l2ctrl, void* stream,
                       uint8_t* x, uint8_t* y, uint8_t* z, uint8_t* out,
                       uint8_t* workspace, uint8_t* tiling) {
  addcdiv_custom<<<blockDim, l2ctrl, stream>>>(x, y, z, out, workspace, tiling);
}
#endif

三、编写自己的算子

1、搭建框架

我们可以使用参考add算子搭建以下目录结构。以下文件夹内的文件没有特别说明就直接从add算子工程内复制。

myCustom
├── AclNNInvocation
│   ├── inc
│   ├── scripts
│   └── src
│   ├── run.sh
├── myCustom 		<-由msopgen工具生成
├── KernelLaunch
│   ├── myCustom.cpp
│   ├── cmake
│   ├── CMakeLists.txt
│   ├── data_utils.h
│   ├── run.sh
│   └── scripts
└── myCustom.json

2、 KernelLaunch编写

1、myCustom.cpp

我们直接cv add算子的,对输入做下修改,然后修改compute就行了。

2、main.cpp

这里主要是将算子名称以及传入的参数修改下
在这里插入图片描述

3、scripts/gen_data.py

这里根据你要实现的代码编写生成数据和真值的程序就行了,在比赛时,我们可以直接从官方给出的AclNN中取。

3、 framework编写

在kernelLaunch测试通过后我们直接修改myCustom.json。如果是多个数据类型,如下所示。

[
    {
        "op": "myCustom",
        "language": "cpp",
        "input_desc": [
            {
                "name": "x",
                "param_type": "required",
                "format": [
                    "ND","ND"
                ],
                "type": [
                    "fp16","fp32"
                ]
            }
        ],
        "output_desc": [
            {
                "name": "y",
                "param_type": "required",
                "format": [
                    "ND","ND"
                ],
                "type": [
                    "fp16","fp32"
                ]
            }
        ]
    }
]

然后生成工程(具体目录请自行修改)

/usr/local/Ascend/ascend-toolkit/latest/python/site-packages/bin/msopgen gen -i /home/HwHiAiUser/myCustom/myCustom.json -c ai_core-ascend310B1 -lan cpp -out /home/HwHiAiUser/myCustom/myCustom

接着就是参考add和addcdiv算子在op_host中实现tiling策略,将kernelLaunch中测试通过的代码加上tiling相关的代码后搬运到op_kernel。编译安装算子。

4、 Aclnn测试

这里因为我做的是比赛里给出的题目,因此直接使用官方给的案例进行测试。对于自定义算子,除修改gen_data外,我们还需要修改op_runner以及main.cpp。

四、torch_npu重新编译(可选)

参考仓库说明:https://gitee.com/ascend/op-plugin

五、常用api

为了简化使用,以下仅列出常用的2级接口,如需高性能实现,请使用0级接口。310b系列似乎不支持高级api,因此也不列出了。详细内容请直接看api文档

名称功能表达式二级接口样例
Exp按元素取自然指数在这里插入图片描述Exp(dstLocal, srcLocal, 512);
Abs按元素取绝对值在这里插入图片描述Abs(dstLocal, srcLocal, 512);
Reciprocal按元素取倒数在这里插入图片描述Reciprocal(dstLocal, srcLocal, 512);
Sqrt按元素做开方在这里插入图片描述Sqrt(dstLocal, srcLocal, 512);
Ln按元素取自然对数在这里插入图片描述Ln(dstLocal, srcLocal, 512);
Add按元素求和在这里插入图片描述Add(dstLocal, src0Local, src1Local, 512);
Mul按元素求积在这里插入图片描述Mul(dstLocal, src0Local, src1Local, 512);
Adds/Muls矢量内每个element与标量求和/积同上Adds(dstLocal, srcLocal, half(2), 512);
Sub按元素求差在这里插入图片描述Sub(dstLocal, src0Local, src1Local, 512);
Div按element求商在这里插入图片描述Div(dstLocal, src0Local, src1Local, 512);
Max按element求最大值在这里插入图片描述Max(dstLocal, src0Local, src1Local, 512);
Min按element求最小值在这里插入图片描述Min(dstLocal, src0Local, src1Local, 512);
Duplicate将一个变量或一个立即数,复制多次并填充到向量在这里插入图片描述Duplicate(dstLocal, half(18.0), 512);

注意:标量双目指令中没有减法和除法,基础api没有log只有ln。

问题

一句话,多看文档,有问题就先去社区搜一下。160001,error code 0这种就直接查代码吧,没有具体原因。

1、fatal error: register/tilingdata_base.h: No such file or directory

在这里插入图片描述
在这里插入图片描述
检查一下CANN路径

其他能稳定复现的bug等我后面遇到了再补充解决办法吧。

总结

也许,有时歪门邪道比正道更简单。不要被文档和案例限制了,不要问能不能,跑下试试最快
就像adds直接乘标量不好使,那就直接把这个标量填满整个local,直接用张量去计算嘛。而且这样能用的api还更多呢。

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mfbz.cn/a/591848.html

如若内容造成侵权/违法违规/事实不符,请联系我们进行投诉反馈qq邮箱809451989@qq.com,一经查实,立即删除!

相关文章

JDK14特性

JDK14 1 概述2 语法层面的变化1_instanceof的模式匹配(预览)2_switch表达式(标准)3_文本块改进(第二次预览)4_Records 记录类型(预览 JEP359) 3 API层面的变化4 关于GC1_G1的NUMA内存分配优化2_弃用SerialCMS,ParNewSerial Old3_删除CMS4_ZGC on macOS and Windows 4 其他变化1…

PPT基础

5种ppt仅可读形式 Ⅰ 开始选项卡 1.【幻灯片】组中&#xff1a;新建幻灯片&#xff0c;从大纲中导入幻灯片&#xff1b;修改幻灯片的版式&#xff1b;节&#xff08;新增节&#xff0c;重命名节&#xff09;。 2.【字体】组中&#xff1a;设置字体&#xff0c;字体大小&…

ctfshow web入门 sql注入 web224--web233

web224 扫描后台&#xff0c;发现robots.txt&#xff0c;访问发现/pwdreset.php &#xff0c;再访问可以重置密码 &#xff0c;登录之后发现上传文件 检查发现没有限制诶 上传txt,png,zip发现文件错误了 后面知道群里有个文件能上传 <? _$GET[1]_?>就是0x3c3f3d60245…

海外仓系统与跨境电商平台集成:有什么意义,为什么重要

跨境电商的发展趋势并没有丝毫放缓的迹象&#xff0c;这使得对高效率、综合性的海外仓的需求变得比以往任何时间都要多。 预测表明&#xff0c;未来一年跨境电商的市场份额将继续扩大。这一切都要求海外仓企业尽快提升仓储管理效率&#xff0c;在这个过程中&#xff0c;海外仓系…

小苹果

题目描述 小的桌子上放着几个苹果从左到右排成一列&#xff0c;编号为从1 到 。小苞是小的好朋友&#xff0c;每天她都会从中拿走一些苹果。每天在拿的时候&#xff0c;小苞都是从左侧第1个苹果开始、每隔2个苹果拿走1个苹果。随后小苞会将剩下的苹果按原先的顺序重新排成一列…

扩展学习|本体研究进展

文献来源&#xff1a; 王向前,张宝隆,李慧宗.本体研究综述[J].情报杂志,2016,35(06):163-170. 一、本体的定义 本体概念被引入人工智能、知识工程等领域后被赋予了新的含义。然而不同的专家学者对本体的理解不同,所给出的定义也有所差异。 人工智能领域的学者Neches(1991)等人对…

StampedLock(戳记锁)源码解读与使用

&#x1f3f7;️个人主页&#xff1a;牵着猫散步的鼠鼠 &#x1f3f7;️系列专栏&#xff1a;Java源码解读-专栏 &#x1f3f7;️个人学习笔记&#xff0c;若有缺误&#xff0c;欢迎评论区指正 1. 前言 我们在上一篇写ReentrantReadWriteLock读写锁的末尾留了一个小坑&#…

这书不错,古琴乐理实用教程(尹溧新编),有课学得通透。

通篇阅读后&#xff0c;发现这本书以古琴初习者、未系统接触过现代乐理的读者为对象&#xff0c;将复杂的古琴音乐理论简单化、通俗化。书中采用参照比较的方法、通俗易懂的语言、言简意赅的文字&#xff0c;并结合具体音乐作品将古琴研习中最主要的、最核心的理论知识进行简明…

进程间通信(3)信号量初识

我最近开了几个专栏&#xff0c;诚信互三&#xff01; > |||《算法专栏》&#xff1a;&#xff1a;刷题教程来自网站《代码随想录》。||| > |||《C专栏》&#xff1a;&#xff1a;记录我学习C的经历&#xff0c;看完你一定会有收获。||| > |||《Linux专栏》&#xff1…

72.网络游戏逆向分析与漏洞攻防-角色与怪物信息的更新-完善利用角色与怪物创建的功能

免责声明&#xff1a;内容仅供学习参考&#xff0c;请合法利用知识&#xff0c;禁止进行违法犯罪活动&#xff01; 如果看不懂、不知道现在做的什么&#xff0c;那就跟着做完看效果 现在的代码都是依据数据包来写的&#xff0c;如果看不懂代码&#xff0c;就说明没看懂数据包…

开源模型应用落地-LangChain高阶-Tools工具-集成agents(四)

一、前言 LangChain 的 tools 是一系列关键组件&#xff0c;它们提供了与外部世界进行交互的能力。通过适当的使用这些组件&#xff0c;可以简单实现如执行网络搜索以获取最新信息、调用特定的 API 来获取数据或执行特定的操作、与数据库进行交互以获取存储的信息等需求。 本章…

MATLAB - 自定义惯性矩阵

系列文章目录 前言 一、关键惯性约定 Simscape 多体软件在惯性定义中采用了一系列约定。请注意这些约定&#xff0c;因为如果手动进行惯性计算&#xff0c;这些约定可能会影响计算结果。如果您的惯性数据来自 CAD 应用程序或其他第三方软件&#xff0c;这些约定还可能影响到您需…

TranslatePress Pro插件下载:一键国际化,让您的网站走向世界

在全球化的今天&#xff0c;一个多语言的网站是连接不同文化和市场的桥梁。TranslatePress Pro插件&#xff0c;作为一款专为WordPress用户设计的多语言解决方案&#xff0c;以其简便的操作和强大的功能&#xff0c;帮助您的网站跨越语言障碍&#xff0c;吸引全球用户。 [Tran…

vector 的模拟实现

目录 1. vector 的核心框架 2. size 和 capacity 以及 empty 3. reserve 和 push_back 4. insert 5. erase 6. copy constructor 6.1. 第一个版本 6.2. 第二个版本 6.3. 第三个版本 7. operator 7.1. 第一个版本 7.2. 第二个版本 7.3. 第三个版本 8. constructor…

用自然语言即可完全控制用户界面;无需调整的文本至图片生成的ID定制方法;OpenAI构建应用指南

✨ 1: PyWinAssistant 用自然语言即可完全控制用户界面 PyWinAssistant是一个突破性的项目&#xff0c;它基于2023年12月31日发布的技术&#xff0c;代表了首个大型行为模型、开源Windows 10/11人工智能框架。这个框架的主要亮点在于它能够通过利用思维可视化&#xff08;Vis…

Java复习第十九天学习笔记(Cookie、Session登录),附有道云笔记链接

【有道云笔记】十九 4.7 Cookie、Session登录 https://note.youdao.com/s/VwpxfEim 一、会话技术简介 生活中会话 我&#xff1a; 小张&#xff0c;你会跳小苹果码&#xff1f; 小张&#xff1a; 会&#xff0c;怎么了&#xff1f; 我&#xff1a; 公司年会上要表演节目&a…

HTML_CSS学习:常用文本属性

一、文本颜色 相关代码&#xff1a; <!DOCTYPE html> <html lang"en"> <head><meta charset"UTF-8"><title>文本颜色</title><style>div{font-size: 90px;}.atguigu1{color: #238c20;}.atguigu2{color: rgb(2…

AI文章框架分析

大家在文章写作的时候结构难免会有点凌乱&#xff0c;但是自己可能无法发现问题所在&#xff0c;那么有没有一款工具可以帮你自动分析你写的文章框架存在的问题&#xff0c;然后并给你详细的分析报告呢&#xff1f;今天给大家介绍一下文件框架分析助手&#xff01; 使用说明 打…

jQuery Moblie 笔记14 开发跨平台移动设备网页

相关内容&#xff1a;jQuery Moblie基础、操作、移动设备仿真器、jQuery Moblie网页实例、jQuery Moblie的UI组件、…… jQuery推出了一套新的函数库jQuery Mobile&#xff0c;目的是希望能够统一当前移动设备的用户界面(UI)。 移动设备开发应用程序目前大致分为两种&#xff…

大数据分析入门之10分钟掌握GROUP BY语法

前言 书接上回大数据分析入门10分钟快速了解SQL。 本篇将会进一步介绍group by语法。 基本语法 SELECT column_name, aggregate_function(column_name) FROM table_name GROUP BY column_name HAVING condition假设我们有students表&#xff0c;其中有id,grade_number,class…
最新文章