Ascend C Add算子样例代码详解

核函数定义

核函数(Kernel Function)是Ascend C算子设备侧实现的入口。在核函数中,需要为在一个核上执行的代码规定要进行的数据访问和计算操作,当核函数被调用时,多个核都执行相同的核函数代码,具有相同的参数,并行执行。

// 实现核函数
extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{// 初始化算子类,算子类提供算子初始化和核心处理等方法KernelAdd op;// 初始化函数,获取该核函数需要处理的输入输出地址,同时完成必要的内存初始化工作op.Init(x, y, z);// 核心处理函数,完成算子的数据搬运与计算等核心逻辑op.Process();
}// 调用核函数
void add_custom_do(uint32_t blockDim, void* l2ctrl, void* stream, uint8_t* x, uint8_t* y, uint8_t* z)
{add_custom<<<blockDim, l2ctrl, stream>>>(x, y, z);
}

使用__global__函数类型限定符来标识它是一个核函数,可以被<<<…>>>调用;使用__aicore__函数类型限定符来标识该核函数在设备端AI Core上执行:

__global__ __aicore__ void kernel_name(argument list);

编程中使用到的函数可以分为三类:核函数(device侧执行)、host侧执行函数、device侧执行函数(除核函数之外的)。三者的调用关系如下图所示:
● host侧执行函数可以调用同类的host执行函数,也就是通用C/C++编程中的函数调用;也可以通过<<<>>>调用核函数。
● device侧执行函数(除核函数之外的)可以调用同类的device侧执行函数。
● 核函数可以调用device侧执行函数(除核函数之外的)。
在这里插入图片描述

Add算子代码分析

Add算子设计规格

在这里插入图片描述

核函数开发

调用算子类的Init和Process函数

extern "C" __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z)
{KernelAdd op;op.Init(x, y, z);op.Process();
}

矢量编程范式实现算子类

class KernelAdd {
public:
__aicore__ inline KernelAdd(){}
// 初始化函数,完成内存初始化相关操作
__aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z){}
// 核心处理函数,实现算子逻辑,调用私有成员函数CopyIn、Compute、CopyOut完成矢量算子的三级流水操作
__aicore__ inline void Process(){}private:
// 搬入函数,完成CopyIn阶段的处理,被核心Process函数调用
__aicore__ inline void CopyIn(int32_t progress){}
// 计算函数,完成Compute阶段的处理,被核心Process函数调用
__aicore__ inline void Compute(int32_t progress){}
// 搬出函数,完成CopyOut阶段的处理,被核心Process函数调用
__aicore__ inline void CopyOut(int32_t progress){}private:
TPipe pipe;  //Pipe内存管理对象
TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY;  //输入数据Queue队列管理对象,QuePosition为VECIN
TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueZ;  //输出数据Queue队列管理对象,QuePosition为VECOUT
GlobalTensor<half> xGm, yGm, zGm;  //管理输入输出Global Memory内存地址的对象,其中xGm, yGm为输入,zGm为输出
};

内部函数调用关系
在这里插入图片描述
由此可见除了Init函数完成初始化外,Process中完成了对流水任务:“搬入、计算、搬出”的调用,开发者可以重点关注三个流水任务的实现。
核函数运行验证
异构计算架构中,NPU(kernel侧)与CPU(host侧)是协同工作的,完成了kernel侧核函数开发后,即可编写host侧的核函数调用程序,实现从host侧的APP程序调用算子,执行计算过程。
代码整体结构
● 调用算子的应用程序:main.cpp。
● 输入数据和真值数据生成脚本文件:gen_data.py。
● 验证输出数据和真值数据是否一致的验证脚本:verify_result.py。
● 编译cpu侧或npu侧运行的算子的编译工程文件:CMakeLists.txt。
● 编译运行算子的脚本:run.sh。
以调用算子的应用程序的编写为例
NPU侧运行算子的调用程序
在这里插入图片描述
完整代码

// AscendCL初始化
CHECK_ACL(aclInit(nullptr));
// 运行管理资源申请
aclrtContext context;
int32_t deviceId = 0;
CHECK_ACL(aclrtSetDevice(deviceId));
CHECK_ACL(aclrtCreateContext(&context, deviceId));
aclrtStream stream = nullptr;
CHECK_ACL(aclrtCreateStream(&stream));
// 分配Host内存
uint8_t *xHost, *yHost, *zHost;
uint8_t *xDevice, *yDevice, *zDevice;
CHECK_ACL(aclrtMallocHost((void**)(&xHost), inputByteSize));
CHECK_ACL(aclrtMallocHost((void**)(&yHost), inputByteSize));
CHECK_ACL(aclrtMallocHost((void**)(&zHost), outputByteSize));
// 分配Device内存
CHECK_ACL(aclrtMalloc((void**)&xDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void**)&yDevice, inputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
CHECK_ACL(aclrtMalloc((void**)&zDevice, outputByteSize, ACL_MEM_MALLOC_HUGE_FIRST));
// Host内存初始化
ReadFile("./input/input_x.bin", inputByteSize, xHost, inputByteSize);
ReadFile("./input/input_y.bin", inputByteSize, yHost, inputByteSize);
CHECK_ACL(aclrtMemcpy(xDevice, inputByteSize, xHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
CHECK_ACL(aclrtMemcpy(yDevice, inputByteSize, yHost, inputByteSize, ACL_MEMCPY_HOST_TO_DEVICE));
// 用内核调用符<<<>>>调用核函数完成指定的运算,add_custom_do中封装了<<<>>>调用
add_custom_do(blockDim, nullptr, stream, xDevice, yDevice, zDevice);
CHECK_ACL(aclrtSynchronizeStream(stream));
// 将Device上的运算结果拷贝回Host
CHECK_ACL(aclrtMemcpy(zHost, outputByteSize, zDevice, outputByteSize, ACL_MEMCPY_DEVICE_TO_HOST));
WriteFile("./output/output_z.bin", zHost, outputByteSize);
// 释放申请的资源
CHECK_ACL(aclrtFree(xDevice));
CHECK_ACL(aclrtFree(yDevice));
CHECK_ACL(aclrtFree(zDevice));
CHECK_ACL(aclrtFreeHost(xHost));
CHECK_ACL(aclrtFreeHost(yHost));
CHECK_ACL(aclrtFreeHost(zHost));
// AscendCL去初始化
CHECK_ACL(aclrtDestroyStream(stream));
CHECK_ACL(aclrtDestroyContext(context));
CHECK_ACL(aclrtResetDevice(deviceId));
CHECK_ACL(aclFinalize());

编译运行

bash run.sh -r npu -v Ascend910A

输出结果

Scanning dependencies of target add_npu
[ 33%] Building CCE object cmake/npu/CMakeFiles/add_npu.dir/__/__/add_custom.cpp.o
[ 66%] Building CCE object cmake/npu/CMakeFiles/add_npu.dir/__/__/main.cpp.o
[100%] Linking CCE executable ../../../add_npu
[100%] Built target add_npu
INFO: compile op on npu succeed!
INFO: execute op on npu succeed!
test pass

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

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

L55--- 257.二叉树的所有路径(深搜)---Java版

1.题目描述 2.思路 &#xff08;1&#xff09;因为是求二叉树的所有路径 &#xff08;2&#xff09;然后是带固定格式的 所以我们要把每个节点的整数数值换成字符串数值 &#xff08;3&#xff09;首先先考虑根节点&#xff0c;也就是要满足节点不为空 返回递归的形式dfs(根节…

AI通用大模型不及垂直大模型?各有各的好

​​​​​​​AI时代&#xff0c;通用大模型和垂直大模型&#xff0c;两者孰优孰劣&#xff0c;一直众说纷纭。 通用大模型&#xff0c;聚焦基础层&#xff0c;如ChatGPT、百度文心一言&#xff0c;科大讯飞星火大模型等&#xff0c;都归属通用大模型&#xff0c;它们可以解答…

51单片机STC89C52RC——4.1 独立按键(数码管显示按键值)

目录 目录 目的 一&#xff0c;STC单片机模块 二&#xff0c;矩阵按键模块 2.1 针脚定义 ​编辑 2.2 矩阵按键位置 2.3 如何理解按键按下后针脚的高低电平 2.3.1 错误理解1 2.3.2 错误理解2 2.3.3 正确判定按下的是那个按键的逻辑 2.3.4 判定按键按下的依次扫描程…

游戏中插入音效

一、背景音乐 准备&#xff1a;素材音乐 方法&#xff1a; 1、方法1&#xff1a; (1) 将背景音乐 bgAudio 拖放到Hierarchy面板 (2) 选中 bgAudio&#xff0c;勾选开始运行就播放、循环播放。调节音量&#xff08;volume) 2、方法2&#xff1a; (1) Create Empty&#x…

大语言模型-Transformer

目录 1.概述 2.作用 3.诞生背景 4.历史版本 5.优缺点 5.1.优点 5.2.缺点 6.如何使用 7.应用场景 7.1.十大应用场景 7.2.聊天机器人 8.Python示例 9.总结 1.概述 大语言模型-Transformer是一种基于自注意力机制&#xff08;self-attention&#xff09;的深度学习…

算法篇-二叉树

二叉树的遍历 分为前序、中序和后续的遍历&#xff0c;思想就是利用递归。 前序遍历-中左右 代码&#xff1a; public void travelTree(TreeNode node, List<Integer> resulst) {if (node null){return;}// 中resulst.add(node.val);// 左travelTree(node.left, resul…

DN-DETR

可以看到&#xff0c;与 DAB-DETR 相比&#xff0c;最大的差别仍然在 decoder 处&#xff0c;主要是 query 的输入。DN-DETR 认为可以把对 offsets 的学习&#xff0c;看作一种对噪声学习的过程&#xff0c;因此&#xff0c;可以直接在 GT 周围生成一些 noised boxes&#xff0…

【机器学习】transformer框架理论详解和代码实现

Hi~&#xff01;这里是奋斗的小羊&#xff0c;很荣幸您能阅读我的文章&#xff0c;诚请评论指点&#xff0c;欢迎欢迎 ~~ &#x1f4a5;&#x1f4a5;个人主页&#xff1a;奋斗的小羊 &#x1f4a5;&#x1f4a5;所属专栏&#xff1a;C语言 &#x1f680;本系列文章为个人学习…

Tower 使用指南

Tower 使用指南 目录 打开 git 仓库查看分支历史切换分支提交修改推送修改创建标签自动拉取最新代码 打开 git 仓库 File -> Open然后选择项目目录 查看分支历史 切换分支 提交修改 推送修改 创建标签 自动拉取最新代码

【阿里云服务器】【弹性云服务ECS】通过ssh登录远程服务器

一、操作系统 使用Windows11主机上的Ubuntu子系统&#xff0c;如下图所示&#xff1a; 二、云服务器登录方法 需知道&#xff1a;服务器ip地址、登录名和自己设置的登录密码&#xff1a; 上述系统用户名为root&#xff0c;需要在Ubuntu子系统中同样切换至root用户&#xff…

3GPP R18冻结,哪些信息值得关注?

这两天网上到处都是R18冻结的新闻&#xff0c;小枣君也凑个热闹&#xff0c;聊聊这个话题。 首先给小白科普一下&#xff0c;所谓的R18&#xff0c;全称是Release-18。它是通信国际标准组织3GPP&#xff08;第三代合作伙伴计划&#xff09;所推出协议标准的一个版本。 自从1998…

Ubuntuwin11双系统

一、准备工作 win11与ubuntu20.4双系统安装案例教程,先查看引导模式参数不服则不要安装否则会报异常 查看BIOS引导模式 查看磁盘分区格式 下载Ubuntu镜像 所有版本下载地址,我的华为云镜像ubuntu20.4这个版本地址

凯迪仕霸榜全渠道TOP1 全域曝光100亿

618年中狂欢盛典&#xff0c;已正式落下帷幕。智能锁行业领头羊凯迪仕&#xff0c;凭借过硬的科技产品力和品牌势能&#xff0c;在全域流量加持以及传奇大师K70新品强势曝光之下&#xff0c;霸榜天猫、京东、抖音各平台&#xff0c;稳居各类型榜单榜首&#xff0c;继续以行业领…

Python xlrd库:读excel表格

&#x1f49d;&#x1f49d;&#x1f49d;欢迎莅临我的博客&#xff0c;很高兴能够在这里和您见面&#xff01;希望您在这里可以感受到一份轻松愉快的氛围&#xff0c;不仅可以获得有趣的内容和知识&#xff0c;也可以畅所欲言、分享您的想法和见解。 推荐:「stormsha的主页」…

带颜色的3D点云数据发布到ros1中(通过rviz显示)python、C++

ros中发布点云数据xyz以及带颜色的点云数据xyzrgb ros中发布点云数据xyz可以直接用python来做或者C(看个人偏好) ros中发布带颜色的点云数据xyzrgb环境1.新建ROS工作空间2.创建功能包 ros中发布点云数据xyz 可以直接用python来做或者C(看个人偏好) 在这里我们带有颜色的点云数…

TikTok 推出专属AI 内容工具

TikTok最近推出了一款极具实用性的新工具包——TikTok Symphony。它融合了生成式人工智能技术&#xff0c;让内容创作变得更加迅速和便捷。 无论是营销人员还是创作者&#xff0c;都能在TikTok上轻松制作出高质量的内容。Symphony将人类的创造力与AI的高效性完美融合&#xff0…

市值3万亿英伟达的崛起:技术、坚持与市场的力量,厚积薄发的经典案例

在科技领域&#xff0c;英伟达&#xff08;NVIDIA&#xff09;的故事无疑是一个厚积薄发的经典案例。作为一家专注于图形处理单元&#xff08;GPU&#xff09;的公司&#xff0c;英伟达用31年的时间证明了技术的价值、计算的价值和坚持的价值。本文将详细探讨英伟达如何从一家市…

车辆轨迹预测系列 (二):常见数据集介绍

车辆轨迹预测系列 (二)&#xff1a;常见数据集介绍 文章目录 车辆轨迹预测系列 (二)&#xff1a;常见数据集介绍1、NuScenes (2020)&#xff1a;1、下载2、说明 2、Waymo Open Dataset (2020)&#xff1a;1、介绍2、概述3、下载4、教程5、参考 3、Lyft Level 5 (2020)&#xff…

SD-WAN为什么适合小企业

SD-WAN&#xff08;软件定义广域网&#xff09;是一种革新性的网络技术&#xff0c;通过软件智能管理&#xff0c;实现灵活和高效的网络连接。在数字化转型浪潮中&#xff0c;企业对网络稳定性和性能的要求不断提升&#xff0c;SD-WAN因此受到了广泛关注。对于资源有限的小型企…

JAVA NIO(二) Buffer和Channel

一&#xff0c;基本使用 1&#xff0c; 一个Socket连接使用一个Channel来表示&#xff0c;以前直接操作Socket文件描述符来对读写缓冲区操作&#xff0c;比如读数据到用户空间的一个byte数组&#xff0c;NIO中Channel对这个过程作了封装&#xff0c;其中用户空间的byte数组就类…