Ascend C Add算子样例代码详解

2024-06-21 22:52
文章标签 代码 详解 算子 add 样例 ascend

本文主要是介绍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

这篇关于Ascend C Add算子样例代码详解的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！


http://www.chinasem.cn/article/1082567。

相关文章

uniapp接入微信小程序原生代码配置方案（优化版）

uniapp接入微信小程序原生代码配置方案（优化版）

uniapp项目需要把微信小程序原生语法的功能代码嵌套过来，无需把原生代码转换为uniapp，可以配置拷贝的方式集成过来 1、拷贝代码包到src目录 2、vue.config.js中配置原生代码包直接拷贝到编译目录中 3、pages.json中配置分包目录，原生入口组件的路径 4、manifest.json中配置分包，使用原生组件 5、需要把原生代码包里的页面修改成组件的方
阅读更多...
公共筛选组件（二次封装antd）支持代码提示

公共筛选组件（二次封装antd）支持代码提示

如果项目是基于antd组件库为基础搭建，可使用此公共筛选组件 使用到的库 npm i antdnpm i lodash-esnpm i @types/lodash-es -D /components/CommonSearch index.tsx import React from 'react';import { Button, Card, Form } from 'antd'
阅读更多...
17.用300行代码手写初体验Spring V1.0版本

17.用300行代码手写初体验Spring V1.0版本

1.1.课程目标 1、了解看源码最有效的方式，先猜测后验证，不要一开始就去调试代码。 2、浓缩就是精华，用 300行最简洁的代码 提炼Spring的基本设计思想。 3、掌握Spring框架的基本脉络。 1.2.内容定位 1、 具有1年以上的SpringMVC使用经验。 2、 希望深入了解Spring源码的人群，对 Spring有一个整体的宏观感受。 3、 全程手写实现SpringM
阅读更多...
十四、观察者模式与访问者模式详解

十四、观察者模式与访问者模式详解

21.观察者模式 21.1.课程目标 1、 掌握观察者模式和访问者模式的应用场景。 2、 掌握观察者模式在具体业务场景中的应用。 3、 了解访问者模式的双分派。 4、 观察者模式和访问者模式的优、缺点。 21.2.内容定位 1、 有 Swing开发经验的人群更容易理解观察者模式。 2、 访问者模式被称为最复杂的设计模式。 21.3.观察者模式 观 察 者 模 式 ( Obser
阅读更多...
【操作系统】信号Signal超详解|捕捉函数

【操作系统】信号Signal超详解|捕捉函数

🔥博客主页： 我要成为C++领域大神🎥系列专栏：【C++核心编程】 【计算机网络】 【Linux编程】 【操作系统】 ❤️感谢大家点赞👍收藏⭐评论✍️ 本博客致力于知识分享，与更多的人进行学习交流 ​ 如何触发信号 信号是Linux下的经典技术，一般操作系统利用信号杀死违规进程，典型进程干预手段，信号除了杀死进程外也可以挂起进程 kill -l 查看系统支持的信号
阅读更多...
Jitter Injection详解

Jitter Injection详解

一、定义与作用 Jitter Injection，即抖动注入，是一种在通信系统中人为地添加抖动的技术。该技术通过在发送端对数据包进行延迟和抖动调整，以实现对整个通信系统的时延和抖动的控制。其主要作用包括： 改善传输质量：通过调整数据包的时延和抖动，可以有效地降低误码率，提高数据传输的可靠性。均衡网络负载：通过对不同的数据流进行不同程度的抖动注入，可以实现网络资源的合理分配，提高整体传输效率。增
阅读更多...
代码随想录算法训练营：12/60

代码随想录算法训练营：12/60

非科班学习算法day12 | LeetCode150:逆波兰表达式 ，Leetcode239: 滑动窗口最大值  目录 介绍 一、基础概念补充： 1.c++字符串转为数字 1. std::stoi, std::stol, std::stoll, std::stoul, std::stoull（最常用） 2. std::stringstream 3. std::atoi, std
阅读更多...
记录AS混淆代码模板

记录AS混淆代码模板

开启混淆得先在build.gradle文件中把 minifyEnabled false改成true,以及shrinkResources true//去除无用的resource文件 这些是写在proguard-rules.pro文件内的 指定代码的压缩级别 -optimizationpasses 5 包明不混合大小写 -dontusemixedcaseclassnames 不去忽略非公共
阅读更多...
一道经典Python程序样例带你飞速掌握Python的字典和列表

一道经典Python程序样例带你飞速掌握Python的字典和列表

Python中的列表（list）和字典（dict）是两种常用的数据结构，它们在数据组织和存储方面有很大的不同。 列表（List） 列表是Python中的一种有序集合，可以随时添加和删除其中的元素。列表中的元素可以是任何数据类型，包括数字、字符串、其他列表等。列表使用方括号[]表示，元素之间用逗号,分隔。 定义和使用 # 定义一个列表 fruits = ['apple', 'banana
阅读更多...
Steam邮件推送内容有哪些？配置教程详解！

Steam邮件推送内容有哪些？配置教程详解！

Steam邮件推送功能是否安全？如何个性化邮件推送内容？ Steam作为全球最大的数字游戏分发平台之一，不仅提供了海量的游戏资源，还通过邮件推送为用户提供最新的游戏信息、促销活动和个性化推荐。AokSend将详细介绍Steam邮件推送的主要内容。 Steam邮件推送：促销优惠 每当平台举办大型促销活动，如夏季促销、冬季促销、黑色星期五等，用户都会收到邮件通知。这些邮件详细列出了打折游戏、
阅读更多...

Copyright China编程 23002807@qq.com

沪ICP备2023033072号-2