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

举报
昇腾CANN 发表于 2024/06/17 14:33:50 2024/06/17
【摘要】 Ascend C是CANN针对算子开发场景推出的编程语言,原生支持C和C++标准规范,最大化匹配用户开发习惯;通过多层接口抽象、自动并行计算、孪生调试等关键技术,极大提高算子开发效率,助力AI开发者低成本完成算子开发和模型调优部署。全新升级的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,将常用算法逻辑进行封装,可减少重复开发,极大提高开发者的开发效率。

Ascend C开放API举例如下:

 图1.png

使用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计算

图2.png

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

图3.png

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

图4.png

以下是通算融合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主要用于打印标量和字符串信息,辅助开发者了解算子中间计算结果,提升调试效率。

图5.png

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更多信息,欢迎登录昇腾社区使用体验:https://www.hiascend.com/ascend-c

【版权声明】本文为华为云社区用户原创内容,转载时必须标注文章的来源(华为云社区)、文章链接、文章作者等基本信息, 否则作者和本社区有权追究责任。如果您发现本社区中有涉嫌抄袭的内容,欢迎发送邮件进行举报,并提供相关证据,一经查实,本社区将立刻删除涉嫌侵权内容,举报邮箱: cloudbbs@huaweicloud.com
  • 点赞
  • 收藏
  • 关注作者

评论(0

0/1000
抱歉,系统识别当前为高风险访问,暂不支持该操作

全部回复

上滑加载中

设置昵称

在此一键设置昵称,即可参与社区互动!

*长度不超过10个汉字或20个英文字符,设置后3个月内不可修改。

*长度不超过10个汉字或20个英文字符,设置后3个月内不可修改。