Ascend C 2.0新特性详解,支撑大模型融合算子高效开发

近日,昇腾算子编程语言Ascend C发布2.0版本,新增支持通算融合MC²特性,使能大模型场景下通信和计算并行,提高整网运行性能;提供更丰富的API覆盖当前主流的融合算子开发场景,提升开发效率;同时通过算子上板调试等调试能力的增强,可将典型算子功能调试耗时降低至小时级。

Ascend C是CANN针对算子开发场景推出的编程语言,原生支持C和C++标准规范,最大化匹配用户开发习惯;通过多层接口抽象、自动并行计算、孪生调试等关键技术,极大提高算子开发效率,助力AI开发者低成本完成算子开发和模型调优部署。全新升级的Ascend C 2.0版本将进一步贴近用户大模型场景的开发诉求,带来更易用的开发体验和更强大的算子执行性能。

1 提供更丰富的API,算子编程效率提升 

目前,Ascend C 2.0已提供130+基础API和60+高阶API极大丰富了Ascend C编程易用性。Ascend C基础API作为算子开发体系的基石,为昇腾AI处理器的矩阵计算、矢量计算、内存管理等能力提供了开放接口,让开发者可以更灵活的对计算过程进行处理;同时,面对更复杂、多变的计算场景,Ascend C提供Matmul、SoftMax等高阶API,将常用算法逻辑进行封装,可减少重复开发,极大提高开发者的开发效率。

使用Matmul高阶API开发者可以四步完成一个矩阵乘计算操作,免去了复杂的tiling算法设计和数据切分搬运,可以高效的完成矩阵乘计算,同时大幅降低了融合算子的开发难度,支撑大模型融合算子高效开发。

// 1、创建Matmul对象

typedef MatmulType<Tposition::GM, CubeFormat::ND, half> aType;

typedef MatmulType<TPosition::GM, CubeFormat::ND, half> bType;

typedef MatmulType<TPosition::GM, CubeFormat::ND, float> cType;

typedef MatmulType<TPosition::GM,CubeFormat::ND, float> biasType;

Matmul<aType,bType,cType,biasType>

mm.Init(&tiling, &tpipe); // 初始化

// 2、设置左矩阵A、右矩阵B、Bias

mm.SetTensorA(gm_a);  // 设置左矩阵A

mm.SetTensorB(gm_b);  // 设置右矩阵B

mm.SetBias(gm_bias); // 设置Bias

// 3、完成矩阵乘操作

while (mm.Iterate()){

    mm.GetTensorC(ub_c);

    ……}

//mm.IterateAll(gm_c);

// 4、结束矩阵乘操作

mm.End():

2 支持通信计算融合,提升整网运行性能

通算融合MC²(Merged Compute and Communication)主要是将通信和计算的操作过程融合,在通信的同时并行执行计算操作,从而将计算总时间从“T通信+T计算”压缩到“max(T通信,T计算)+拖尾开销(max(t1,t2)/通信轮次)”。以allgather+matmul+add场景举例,支持通算融合MC²特性后,GPT-3的整网性能可提升10%。

通信计算串行方案:计算总时间 = T通信+T计算

MC²通算融合方案:计算总时间 =max(T通信,T计算)+拖尾开销(max(t1,t2)/通信轮次)

Ascend C支持在算子的Kernel代码中进行通信相关操作将集合通信和计算进行深度融合从最底层的算子侧提供通信和计算融合。

以下是通算融合MC²算子的kernel编程伪码:

__global__ __aicore__ void Kernel(GMADDR gmA, gmB, gmC, Tiling* tiling) {

  Matmul<…> mm; //完成Matmul 的定义

  …  

HCCLComm<HCCLDmaMode::SDMA> comm;      

comm.Init(tiling->ptrHCCLTiling); //在Host Tiling时填充,初始化通信对象          

gmWin = comm.GetCommAddr();      //获得通信域地址,减少Copy   

//先通知通信引擎准备数据,并把数据放到win窗口  

comm.Prepare<AllGather>(gmWin, gmA, size, DATA_TYPE_FP16);  

comm.Commit(); //发起一次通信,让通信先发起,然后本地计算  

//计算本地的一片数据做矩阵计算 A*B=C ,此时通信和计算异步并行  

mm.SetTensorA(gmA);

mm.SetTensorB(gmB);  

mm.IterateAll(gmC + offset);     

//gather后数据计算  

for(int i=0; i < tiling.tileCnt; i++) { //对A矩阵Gather数据切分成tileCnt份, 

   if( i != tileCnt-1) { //通知通信引擎,进行下一次通信

       comm.Prepare<AllGather>(gmWin+offsetWin, gmA+offsetA, size, DATA_TYPE_FP16);

      comm.Commit();

    }

    comm.Wait();      //等待Gather数据完成

    SyncAll(); //多核同步,其他核和发起通信的核同步

    mm.SetTensorA(gmWin+offsetWin); //计算接收过来的数据的MM

    mm.IterateAll(gmC+offsetC);     //计算接收过来的数据的MM

  }

}

3 调试能力增强,提高问题定位效率

Ascend C提供了丰富的调试调优手段帮助开发者提高问题定位效率Ascend C 2.0在支持CPU域和NPU域孪生调试的基础上,丰富了NPU域的多种调试方法,可针对性发现和解决各类算子问题。NPU上板调试可针对单算子核函数进行上板运行调试,快速定位算子精度问题;Tiling调试可以进行Tiling函数验证及结果解析,帮助定位Tiling问题;上板Profiling数据采集可在真实环境统计算子执行耗时,帮助快速进行算子性能调优;同时支持指令级流水性能图,方便开发者查看算子内部流水并行处理情况,从而辅助开发者进行性能调试。

其中,NPU域上板调试功能可支持DumpTensor、printf两种数据打印能力,DumpTensor用于打印指定Tensor的数据, printf主要用于打印标量和字符串信息,辅助开发者了解算子中间计算结果,提升调试效率。

 

Ascend C作为昇腾专门面向算子开发场景的编程语言,已经在大模型场景下发挥了巨大的作用,目前,基于Ascend C开发的FlashAttention算子相比融合前的小算子取得了5倍以上的性能提升,开发者可直接调用相关算子API接口使能大模型极致性能优化。相信随着Ascend C 2.0的全新升级,基于Ascend C开发融合算子的优势将愈加凸显,助力开发者突破技术壁垒,为大模型创新提供新方向。

目前,基于Ascend C的开发者大赛和活动正在火热进行中。面向企业的“CANN技术行”、面向高校和普通开发者的“CANN训练营”正在源源不断地为昇腾注入开发源动力,鼓励开发者基于Ascend C的基础能力进行原生开发与实践创新,同时“昇腾AI原生创新算子挑战赛(S2赛季)”也将于近日火热开启。了解Ascend C 2.0更多信息,欢迎登录昇腾社区使用体验:昇腾Ascend C-入门课程-学习资源-算子文档-昇腾社区

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

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

相关文章

大语言模型 (LLM) 红队测试:提前解决模型漏洞

大型语言模型 (LLM) 的兴起具有变革性&#xff0c;以其在自然语言处理和生成方面具有与人类相似的卓越能力&#xff0c;展现出巨大的潜力。然而&#xff0c;LLM 也被发现存在偏见、提供错误信息或幻觉、生成有害内容&#xff0c;甚至进行欺骗行为的情况。一些备受关注的事件包括…

clipboard.js(web页面实现点击复制)

文章目录 codeshow 一个很简单的需求&#xff0c;一个单页面需要一个点击复制的功能 后来在线上找到一个clipboard.js可以实现&#xff0c;这里只用到了最基础的用法&#xff0c;页面样式布局基于bootstrap5.2.3 code <div class"d-flex align-items-center justify-co…

字符设备驱动

目录 demo.c test.c led.h makefile 实验效果 demo.c #include <linux/init.h> #include <linux/module.h> #include <linux/fs.h> #include <linux/uaccess.h> #include <linux/io.h> #include "myled.h" //内核buf char kbuf[…

UITableView之cell复用

关于cell复用的必要性 cellForRowAtIndexPath会随着屏幕滚动而调用&#xff0c;每次出现新行时因为行号变化&#xff0c;就会被调用。 底层原理&#xff1a;当前单元格滚出屏幕时cell销毁&#xff0c;当前单元格又滚回来时cell创建。短时间内频繁创建和销毁cell会影响系统性能…

MySQL 触发器(实验报告)

一、实验名称&#xff1a; 触发器 二、实验日期&#xff1a; 2024 年 6月 8日 三、实验目的&#xff1a; 掌握MySQL触发器的创建及调用&#xff1b; 四、实验用的仪器和材料&#xff1a; 硬件&#xff1a;PC电脑一台&#xff1b; 配置&#xff1a;内存&#xff0c;…

React Native将 ipad 端软件设置为横屏显示后关闭 Modal 弹窗报错

问题&#xff1a; 将 ipad 端软件设置为横屏显示后&#xff0c;关闭 Modal 弹窗报错。 Modal was presented with 0x2 orientations mask but the application only supports 0x18.Add more interface orientations to your apps Info.plist to fix this.NOTE: This will cras…

幸狐RV1106开发板烧录Ubuntu系统与配置SDK,RV1106 LuckFox Pico Max——最新的操作

资料&#xff1a;上手教程 | LUCKFOX WIKI 以及SDK内的文档资料 开发板型号&#xff1a;RV1106 LuckFox Pico Max 烧录系统&#xff1a; Ubuntu 虚拟机系统&#xff1a;Ubuntu 20.04&&Ubuntu22.04 PC系统&#xff1a;win11 占用空间&#xff1a;大概15G 本文主要记…

基于jeecgboot-vue3的Flowable流程-流程处理(一)

因为这个项目license问题无法开源&#xff0c;更多技术支持与服务请加入我的知识星球。 这部分修正一些流程处理中VForm3线上的一些bug问题 1、初始化流程提交与现实的前端页面代码 <!--初始化流程加载默认VForm3表单信息--><el-col :span"16" :offset&qu…

在 Selenium 中更改 User-Agent | 步骤与最佳实践

在 Selenium 中更改 User Agent 是许多网页抓取任务中的关键步骤。它有助于将自动化脚本伪装成常规浏览器&#xff0c;从而避免被网站检测到。本指南将带您了解如何在 Selenium 中更改 Google Chrome 的 User Agent&#xff0c;并提供最佳实践以确保您的网页抓取任务顺利进行。…

IO多路复用简介和代码实例(select函数)

接上篇 阻塞IO、非阻塞IO、IO多路复用和信号驱动IO简介-CSDN博客文章浏览阅读95次。阻塞IO、非阻塞IO、IO多路复用和信号驱动IO简介https://blog.csdn.net/CSDN_DU666666/article/details/139598410?csdn_share_tail%7B%22type%22%3A%22blog%22%2C%22rType%22%3A%22article%2…

【Linux】易错点——/etc/passwd ; /etc/shadow;ifconfig;route;chmod;ps;mv

/etc/passwd ; /etc/shadow /etc/passwd&#xff1a; 用户账户的详细信息在此文件中更新。 用户名&#xff1a;密码&#xff1a;用户 ID&#xff1a;群组 ID&#xff1a;用户 ID 信息&#xff1a;用户的家目录&#xff1a; Shell /etc/shadow&#xff1a; 用户账户密码在此文…

ssm160基于Java技术的会员制度管理的商品营销系统的设计与实现+vue

商品营销系统计与实现 摘 要 现代经济快节奏发展以及不断完善升级的信息化技术&#xff0c;让传统数据信息的管理升级为软件存储&#xff0c;归纳&#xff0c;集中处理数据信息的管理方式。本商品营销系统就是在这样的大环境下诞生&#xff0c;其可以帮助管理者在短时间内处理…

【C++取经之路】继承

目录 继承的概念及定义 单继承的格式 继承方式和访问限定符 继承后子类访问基类成员的权限 基类和派生类对象赋值转换 切片 继承中的作用域 引申&#xff1a;重载和隐藏的区别 派生类的默认成员函数 继承与友元 继承与静态成员 如何实现一个不能被继承的类 复杂的…

【SkiaSharp绘图03】SKPaint详解(一)BlendMode混合模式、ColorFilter颜色滤镜

文章目录 SKPaintSKPaint属性BlendMode获取或设置混合模式SKBlendMode 枚举成员效果预览 Color/ColorF获取或设置前景色ColorFilter 颜色滤镜CreateBlendMode 混合模式CreateColorMatrix 颜色转换CreateCompose 组合滤镜CreateHighContrast 高对比度滤镜CreateLighting 照明滤镜…

Java最新面试题(全网最全、最细、附答案)

一、Java基础 1、基础概念与常识Java 语言有哪些特点? 简单易学&#xff08;语法简单&#xff0c;上手容易&#xff09;&#xff1b;面向对象&#xff08;封装&#xff0c;继承&#xff0c;多态&#xff09;&#xff1b;平台无关性&#xff08; Java 虚拟机实现平台无关性&a…

千万级流量冲击下,如何保证极致性能

1 简要介绍 随着互联网的快速发展&#xff0c;网络应用的流量规模不断攀升&#xff0c;特别是在电商大促、明星直播、重大赛事、头条热搜等热点事件中&#xff0c;秒级100w请求成为了常态。在这样的流量冲击下&#xff0c;如何确保系统稳定、高效地处理每一个请求&#xff0c;为…

抖某音号解封释放实名

##抖音账号封禁后如何解封呢 我相信&#xff0c;做过抖音&#xff0c;或者正在做抖音的朋友&#xff0c;都曾面临一种尴尬至极的局面&#xff0c;辛辛苦苦做起来的账号&#xff0c;或者刚刚准备好的账号&#xff0c;在一时之间&#xff0c;竟然被抖音官方封禁了&#xff01; 实…

ubuntu下使用cmake编译opencv4.8.0+ffmpeg4.2.2+cuda11.1

1.源码下载 &#xff08;1&#xff09;下载ffmpeg4.2.2、opencv4.8.0源码&#xff0c;这里提供一个百度网盘地址&#xff1a; 链接&#xff1a;https://pan.baidu.com/s/1pBksr0_RtKL0cM6Gsf2MGA?pwdcyai 提取码&#xff1a;cyai &#xff08;2&#xff09;解压所有文件 例…

小而美的算法技巧:前缀和数组

小而美的算法技巧&#xff1a;前缀和数组 类似动态规划。 class NumArray {private int[] preSum;public NumArray(int[] nums) {preSumnew int[nums.length1];//preSum[0]的前缀和为0for(int i1;i<preSum.length;i){preSum[i]nums[i-1]preSum[i-1];//先计算累加和}}publi…

Git进阶使用(图文详解)

文章目录 Git概述Git基础指令Git进阶使用一、Git分支1.主干分支2.其他分支2.1创建分支2.2查看分支1. 查看本地分支2. 查看远程分支3. 查看本地和远程分支4. 显示分支的详细信息5. 查看已合并和未合并的分支 2.3切换分支1. 切换到已有的本地分支2. 创建并切换到新分支3. 切换到远…
最新文章