算子迁移指南【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion函数部分函数主体cube/vector和函数部分分别使用__DAV_CUBE__和__DAV_VEC__宏进行隔离函数名部分原有__global__标识替换为__spk__标识函数名更新为其他具备判别性的函数名函数参数部分统一修正为使用__gm__ uint64_t *param指针进行传递通过(GM_ADDR)param[index]方式获取原有参数结构体参数需要通过GET_STRUCT_PTR宏进行获取注意获取到的内容为结构体指针功能主体部分原有逻辑保持不变注意对应的结构体部分转换为了tiling指针无需再次取地址函数结尾不需要再添加pipe_barrier(PIPE_ALL);由框架统一处理额外注意【Planing】目前纯V算子为保证直调使用暂时使用内部方案FunLevelKType进行标识后续会有更完善的方案【后续一定会删除】mix1:1和mix1:2算子需要区分mix_aic和mix_aiv两种类型分别对应AIC和AIV的场景【Planing】目前SK算子暂时没有使用模板写法主要目前暂不能自动获取对应SK入口的函数名后续会完善该部分【Doing】目前8.3版本的cann包条件下不支持在sk框架中使用AscendC::printf, 后续的版本已定位并解决了这个问题编译部分使用ASC编译方式进行编译需要增加ops的common路径内置了迁移所需的辅助宏SK框架部分【Planing】目前在迁移完成后需要在sk_task.cpp的g_sk_fun_map中手动增加函数名映射以便框架能够正确获取函数地址SkBuildTask构建算子任务时当前框架自动偏移了ffts但后续的cann版本底层接口已完成偏移无需框架完成暂时为了兼容仍保留的该部分代码后续会删除该部分逻辑算子实例算子实现编译同时写2个入口函数一个global应用于单算子一个aicore用于sk调用// 默认打开了SuperKernel选项但算子模式退出前自己加一个PIPE_ALL templatetypename DTYPE_X, typename DTYPE_GAMMA __global__ __aicore__ void rms_norm(GM_ADDR x, GM_ADDR gamma, GM_ADDR y, GM_ADDR rstd, RMSNormTilingData tiling) { { KernelRmsNormMergeNDTYPE_X, DTYPE_GAMMA, RMSNormTilingData op; op.Init(x, gamma, y, rstd, tiling); op.Process(); } pipe_barrier(PIPE_ALL); } // aicore函数需要通过宏隔离一下否则编译有问题(当前有bug会导致编译1个aic1个aiv) #ifdef __DAV_VEC__ extern C __aicore__ void rms_norm_half_half_sk(__gm__ uint64_t *param) { GM_ADDR x (GM_ADDR)param[0]; // 参数需要额外处理一下从单一的buffer中获取 GM_ADDR gamma (GM_ADDR)param[1]; GM_ADDR y (GM_ADDR)param[2]; GM_ADDR rstd (GM_ADDR)param[3]; KernelRmsNormMergeNhalf, half, __gm__ RMSNormTilingData op; // tiling结构体需要带__gm__描述不属于参数表了 __gm__ RMSNormTilingData *tilingdata (__gm__ RMSNormTilingData *)(__gm__ uint8_t *)(param 4); op.Init(x, gamma, y, rstd, tilingdata); op.Process(); } #endif【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考