3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

这篇具有很好参考价值的文章主要介绍了3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例。希望对大家有所帮助。如果存在错误或未考虑完全的地方,请大家不吝赐教,您也可以点击"举报违法"按钮提交疑问。

本文分享自华为云社区《3天上手Ascend C编程 | Day2 通过Ascend C编程范式实现一个算子实例》,作者:昇腾CANN 。

一、Ascend C编程范式

Ascend C编程范式把算子内部的处理程序,分成多个流水任务( stage ),以张量( Tensor)为数据载体,以队列 ( Queue ) 进行任务之间的通信与同步,以内存管理模块( Pipe ) 管理任务间的通信内存。

1、流水任务

流水任务指的是单核处理程序中主程序调度的并行任务。在核函数内部,可以通过流水任务实现数据的并行处理,进一步提升性能。下面举例来说明,流水任务如何进行并行调度。以下面的流水任务示意图为例,单核处理程序的功能被拆分成3个流水任务:Stage1、Stage2、Stage3,每个任务专注于完成单一功能;需要处理的数据被切分成n片,使用Progress1~n表示,每个任务需要依次完成n个数据切片的处理。Stage间的箭头表达数据间的依赖关系,比如Stage1处理完Progress1之后,Stage2才能对Progress1进行处理。

 文章来源地址https://www.toymoban.com/news/detail-705958.html

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

若n=3,即待处理的数据被切分成3片,则上图中的流水任务运行起来的示意图如下,从运行图中可以看出,对于同一片数据,Stage1、Stage2、Stage3之间的处理具有依赖关系,需要串行处理;不同的数据切片,同一时间点,可以有多个任务在并行处理,由此达到任务并行、提升性能的目的。

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

 

矢量(Vector)编程范式把算子的实现流程分为3个基本任务:CopyIn,Compute,CopyOut。CopyIn负责搬入操作,Compute负责矢量计算操作,CopyOut负责搬出操作。

 

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

 

2、任务间通信与同步

不同的流水任务之间存在数据依赖,需要进行数据传递。Ascend C中使用Queue队列完成任务之间的数据通信和同步,提供EnQue、DeQue等基础API。Queue队列管理NPU上不同层级的物理内存时,用一种抽象的逻辑位置(QuePosition)来表达各级别的存储,代替了片上物理存储的概念,开发者无需感知硬件架构。

矢量编程中使用到的逻辑位置(QuePosition)定义如下:

搬入数据的存放位置:VECIN搬出数据的存放位置:VECOUT矢量编程主要分为CopyIn、Compute、CopyOut三个任务:

CopyIn任务中将输入数据从Global内存搬运至Local内存后,需要使用EnQue将LocalTensor放入VECIN的Queue中;

Compute任务等待VECIN的Queue中LocalTensor出队之后才可以完成矢量计算,计算完成后使用EnQue将计算结果LocalTensor放入到VECOUT的Queue中;

CopyOut任务等待VECOUT的Queue中LocalTensor出队,再将其拷贝到Global内存。

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

 

Ascend C使用GlobalTensor和LocalTensor作为数据的基本操作单元,它是各种指令API直接调用的对象,也是数据的载体。

3、内存管理机制

任务间数据传递使用到的内存统一由内存管理模块Pipe进行管理。Pipe作为片上内存管理者,通过InitBuffer接口对外提供Queue内存初始化功能,开发者可以通过该接口为指定的Queue分配内存。

Queue队列内存初始化完成后,需要使用内存时,通过调用AllocTensor来为LocalTensor分配内存,当创建的LocalTensor完成相关计算无需再使用时,再调用FreeTensor来回收LocalTensor的内存。

 

编程过程中使用到的临时变量内存同样通过Pipe进行管理。临时变量可以使用TBuf数据结构来申请指定QuePosition上的存储空间,并使用Get()来将分配到的存储空间分配给新的LocaLTensor从TBuf上获取全部长度,或者获取指定长度的LocalTensor。

 

使用TBuf申请的内存空间只能参与计算,无法执行Queue队列的入队出队操作。

二、使用Ascend C编程范式实现一个算子实例

矢量算子开发一般开发流程如下:

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

 

下面以add作为例子介绍Ascend C矢量算子的开发流程。完整样例大家可以参考代码样例。

1、算子分析

分析算子的数学表达式、输入、输出以及计算逻辑的实现,明确需要调用的Ascend C接口。

例子以Add算子为例,数学公式:z= x+y,为简单起见,设定输入张量x, y,z为固定shape(8,2048),数据类型dtype为half类型,数据排布类型format为ND,核函数名称为add_custom。

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

 

 

  • 算子的数学表达式及计算逻辑。Add算子的数学表达式为:z = x + y;计算逻辑:输入数据需要先搬入到片上存储,然后使用计算接口完成两个加法运算,得到最终结果,再搬出到外部存储。

  • 输入和输出。Add算子有两个输入:x与y,输出为z。输入数据类型为half,输出数据类型与输入数据类型相同。输入支持固定shape(8,2048)输出shape与输入shape相同,输入数据排布类型为ND。

  • 确定核函数名称和参数。自定义核函数名,如add_custom。根据输入输出,确定核函数有3个入参x,y,z。x,y为输入在Global Memory上的内存地址,z为输出在Global Memory上的内存地址。

  • 确定算子实现所需接口。涉及内外部存储间的数据搬运,使用数据搬移接口:Datacopy实现;涉及矢量计算的加法操作,使用矢量双目指令:Add实现;使用到LocalTensor,使用Queue队列管理,会使用到EnQue、DeQue等接口。

2、核函数定义

在add_custom核函数的实现中实例化kernelAdd算子类,调用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();
}

3、根据矢量编程范式实现算子类

根据前面的知识,算子实现三个流水任务CopyIn、Compute、CopyOut。任务间通过队列VECIN、VECOUT进行通信和同步,由pipe内存管理对象对任务间交互使用到的内存、临时变量使用到的内存统一进行管理。如下图所示:

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

 

  • CopyIn任务:将Global Memory上的输入Tensor xGm和yGm搬运至Local Memory,分别存储在xLocal,yLocal;

  • Compute任务:对xLocal,yLocal执行加法操作,计算结果存储在zLocal中;

  • CopyOut任务:将输出数据从zLocal搬运至Global Memory上的输出Tensor zGm中

CopyIn,Compute任务间通过VECIN队列inQueuex,inQueuer进行通信和同步;compute,copyout任务间通过VECOUT队列outQueuez进行通信和同步。

第一步,我们进行算子类定义:​

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

 

 

 

 

​第二步,实现Init()函数:多核并行+单核处理数据:

 

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

第三步,实现Process()函数——CopyIn,Compute、CopyOut三个流水任务:

 

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

 

​第四步,通过double buffer机制进一步提升性能

double buffer通过将数据搬运与矢量计算并行执行以隐藏数据搬运时间并降低矢量指令的等待时间,最终提高矢量计算单元的利用效率1个Tensor同一时间只能进行搬入、计算和搬出三个流水任务中的一个,其他两个流水任务涉及的硬件单元则处于ldle状态如果将待处理的数据一分为二,比如Tensor1、Tensor2:

当矢量计算单元对Tensor1进行Compute时,Tensor2可以执行CopvIn的任务

当矢量计算单元对Tensor2进行Compute时,Tensor1可以执行CopyOut的任务

当矢量计算单元对Tensor2进行CopyOut时,Tensor1可以执行CopyIn的任务。由此,数据的进出搬运和矢量计算之间实现并行,硬件单元闲置问题得以有效缓解​

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

 

​最后,基于内核调用符方式进行算子验证

先使用python脚本生成x,y,并计算出z(golden)并落盘。然后再用相同的x,y,在cpu和npu模式下调用add算子,计算出结果z,并与python脚本采用计算md5sum的方式进行对比,完全一样,则表示结果正确。

为了运行方便,我们使用一个run.sh,写有cpu和npu模式的编译命令,通过输入参数进行选择cpu或npu模式进行编译,运行。

1)CPU模式下:

使用ICPU_RUN_KF宏进行CPU调试。

ICPU_RUN_KF(add_custom, blockDim, x, y, z); // use this macro for cpu debug

bash run.sh add_custom ascend910 AiCore cpu运行结果:​

3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例

2)NPU模式下:

NPU模式使用<<<>>>方式调用,由于CPU模式g++没有<<<>>>的表达,需要使用内置宏 __CCE_KT_TEST。

#ifndef __CCE_KT_TEST__
//call of kernel function
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);
}
#endif

bash run.sh add_custom ascend910 AiCore npu运行结果:​

​更多学习资源

好啦,本次分享结束啦,Ascend C的学习资源还有很多,想深入学习的可以参考官网教程:Ascend C官方教程

3天上手Ascend C编程 | Day1 Ascend C基本概念及常用接口3天上手Ascend C编程 | Day2 通过Ascend C编程范式实现一个算子实例3天上手Ascend C编程 | Day3 Ascend C算子调试调优方法

点击关注,第一时间了解华为云新鲜技术~

 

到了这里,关于3天上手Ascend C编程丨通过Ascend C编程范式实现一个算子实例的文章就介绍完了。如果您还想了解更多内容,请在右上角搜索TOY模板网以前的文章或继续浏览下面的相关文章,希望大家以后多多支持TOY模板网!

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处: 如若内容造成侵权/违法违规/事实不符,请点击违法举报进行投诉反馈,一经查实,立即删除!

领支付宝红包 赞助服务器费用

相关文章

  • Ascend C 自定义算子 Kernel Launch调用入门

    本文分享自华为云社区《Ascend C 自定义算子 Kernel Launch调用入门》,作者: jackwangcumt。 根据官方说明文档的介绍,Ascend C对外开放核函数的基础调用(Kernel Launch)方式,是为了简化Ascend C 自定义算子的开发流程,提供更易用的调试调优功能。当开发者完成算子核函数的开发和

    2024年04月09日
    浏览(61)
  • Socket实例,实现多个客户端连接同一个服务端代码&TCP网络编程 ServerSocket和Socket实现多客户端聊天

    Java socket(套接字)通常也称作\\\"套接字\\\",用于描述ip地址和端口,是一个通信链的句柄。应用程序通常通过\\\"套接字\\\"向网络发出请求或者应答网络请求。 使用socket实现多个客户端和同一客户端通讯;首先客户端连接服务端发送一条消息,服务端接收到消息后进行处理,完成后再

    2024年02月12日
    浏览(70)
  • 零编程经验,通过 GPT-4 十分钟开发了一个浏览器插件,并成功运行,实现了需求目标!

    大佬蓝鸟ID: sundyme 零编程经验,通过 GPT-4 十分钟开发了一个浏览器插件,并成功运行,实现了需求目标!太不可思意了,真正体会到了自然语言编程的魅力! 下一步是利用Pinterest 的 API 接口实现自动发图,已经生成好了代码和步骤(看着挺靠谱),等明天开发者权限审核下

    2023年04月08日
    浏览(47)
  • 【RabbitMQ上手——单实例安装&5种简单模式实现通讯过程】

    安装环境:虚拟机VMWare + Centos7.6 + Maven3.6.3 + JDK1.8 RabbitMQ版本:rabbitmq-server-3.8.8-1.el7.noarch.rpm 具体安装过程,可参考:CentOS7安装RabbitMQ(rpm包方式) 用户角色创建 RabbitMQ在安装好后,可以访问http://localhost:15672 ;其自带了guest/guest的用户名和密码;如果需要创建自定义用户;那么

    2024年02月13日
    浏览(35)
  • 纯干货!一文get昇腾Ascend C编程入门全部知识点

    本文分享自华为云社区《昇腾Ascend C编程入门教程》,作者:昇腾CANN 。 2023年5月6日,在昇腾AI开发者峰会上,华为正式发布了面向算子开发场景的昇腾Ascend C编程语言。Ascend C原生支持C/C++编程规范,通过多层接口抽象、并行编程范式、孪生调试等技术,极大提高了算子的开发

    2024年02月10日
    浏览(35)
  • 探索 JavaScript ES8 中的函数式编程并通过实例加以实践

    💂 个人网站:【 海拥】【神级代码资源网站】【办公神器】 🤟 基于Web端打造的:👉轻量化工具创作平台 💅 想寻找共同学习交流的小伙伴,请点击【全栈技术交流群】 函数式编程是一种强大的范式,强调使用纯函数和不可变数据。在本文中,我们将通过实际示例探讨如何

    2024年02月22日
    浏览(45)
  • 【设计模式与范式:行为型】71 | 命令模式:如何利用命令模式实现一个手游后端架构?

    设计模式模块已经接近尾声了,现在我们只剩下 3 个模式还没有学习,它们分别是:命令模式、解释器模式、中介模式。这 3 个模式使用频率低、理解难度大,只在非常特定的应用场景下才会用到,所以,不是我们学习的重点,你只需要稍微了解,见了能认识就可以了。 今天

    2024年02月09日
    浏览(39)
  • 【设计模式与范式:行为型】61 | 策略模式(下):如何实现一个支持给不同大小文件排序的小程序?

    上一节课,我们主要介绍了策略模式的原理和实现,以及如何利用策略模式来移除 if-else 或者 switch-case 分支判断逻辑。今天,我们结合“给文件排序”这样一个具体的例子,来详细讲一讲策略模式的设计意图和应用场景。 除此之外,在今天的讲解中,我还会通过一步一步地

    2024年02月10日
    浏览(40)
  • Flink的常用算子以及实例

    特性:接收一个数据,经过处理之后,就返回一个数据 1.1. 源码分析 我们来看看map的源码 map需要接收一个MapFunctionT,R的对象,其中泛型T表示传入的数据类型,R表示经过处理之后输出的数据类型 我们继续往下点,看看MapFunctionT,R的源码 这是一个接口,那么在代码中,我们就需

    2024年02月12日
    浏览(41)
  • 【云原生-K8s-1实例】通过yaml 文件编排一个web-MySQL小项目

    🍁 博主简介   🏅云计算领域优质创作者   🏅华为云开发者社区专家博主   🏅阿里云开发者社区专家博主 💊 交流社区: 运维交流社区 欢迎大家的加入! RC(ReplicationController)是Kubernetes系统中的核心概念之一,简单来说,它其实定义了一个期望的场景,即声明某

    2024年02月10日
    浏览(62)

觉得文章有用就打赏一下文章作者

支付宝扫一扫打赏

博客赞助

微信扫一扫打赏

请作者喝杯咖啡吧~博客赞助

支付宝扫一扫领取红包,优惠每天领

二维码1

领取红包

二维码2

领红包